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 buildI32Constant(uint32_t Val, MachineInstr &
I,
489 SPIRVTypeInst ResType =
nullptr)
const;
490 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
491 SPIRVTypeInst ResType =
nullptr)
const;
493 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
494 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
495 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
497 MachineInstr &
I)
const;
498 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
500 MachineInstr &
I)
const;
502 bool wrapIntoSpecConstantOp(MachineInstr &
I,
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)
613 LLT DstType = MRI.
getType(DstReg);
615 LLT SrcType = MRI.
getType(SrcReg);
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()) {
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: {
1114 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
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) &&
1293 getImm(
I.getOperand(2), MRI));
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);
1630 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1637 const bool IsUntyped =
1638 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1640 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1641 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1642 : SPIRV::OpVariable))
1645 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1651 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1654 .
addImm(
static_cast<uint32_t
>(Ex.first))
1656 .
add(
I.getOperand(2))
1660 Register ExpResReg =
I.getOperand(1).getReg();
1662 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1672bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1673 SPIRVTypeInst ResType,
1674 MachineInstr &
I)
const {
1675 Register XReg =
I.getOperand(1).getReg();
1676 Register ExpReg =
I.getOperand(2).getReg();
1682 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1683 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1685 ExpReg = buildVectorSplat(ExpReg, NumElts,
I);
1688 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1689 true,
false, {XReg, ExpReg});
1692bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1693 SPIRVTypeInst ResType,
1694 MachineInstr &
I)
const {
1695 Register CosResVReg =
I.getOperand(1).getReg();
1696 unsigned SrcIdx =
I.getNumExplicitDefs();
1701 MachineIRBuilder MIRBuilder(
I);
1703 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1710 const bool IsUntyped =
1711 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1713 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1714 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1715 : SPIRV::OpVariable))
1718 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1722 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1725 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1727 .
add(
I.getOperand(SrcIdx))
1731 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1739 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1742 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1744 .
add(
I.getOperand(SrcIdx))
1746 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1749 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1751 .
add(
I.getOperand(SrcIdx))
1758bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1759 SPIRVTypeInst ResType,
1762 unsigned Opcode)
const {
1763 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1773bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1774 SPIRVTypeInst ResType,
1777 unsigned Opcode)
const {
1778 MachineIRBuilder MIRBuilder(
I);
1780 Register OpReg =
I.getOperand(1).getReg();
1786 SPIRVTypeInst ExtType =
1793 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1797 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1800 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1803bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1804 SPIRVTypeInst ResType,
1807 unsigned Opcode)
const {
1808 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1811bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1812 SPIRVTypeInst ResType,
1815 unsigned Opcode)
const {
1816 MachineIRBuilder MIRBuilder(
I);
1825 SPIRVTypeInst WorkingType =
1832 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {SrcReg}, SPIRV::OpUConvert))
1836 if (!selectOpWithSrcs(LowCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1844 IsScalar ? SPIRV::OpShiftRightLogicalS : SPIRV::OpShiftRightLogicalV;
1846 if (!selectOpWithSrcs(Shift, SrcType,
I, {SrcReg, ShiftAmount}, ShiftOp))
1850 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {Shift}, SPIRV::OpUConvert))
1854 if (!selectOpWithSrcs(HighCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1859 if (!selectOpWithSrcs(Sum, WorkingType,
I, {HighCount, LowCount},
1860 IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV))
1864 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1865 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1868bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1869 SPIRVTypeInst ResType,
1871 unsigned Opcode)
const {
1876 if (!STI.getTargetTriple().isVulkanOS())
1877 return selectUnOp(ResVReg, ResType,
I, Opcode);
1879 Register OpReg =
I.getOperand(1).getReg();
1882 : SPIRV::OpUConvert;
1886 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1888 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1890 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1892 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1896bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1897 SPIRVTypeInst ResType,
1899 unsigned Opcode)
const {
1901 Register SrcReg =
I.getOperand(1).getReg();
1906 unsigned DefOpCode = DefIt->getOpcode();
1907 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1910 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1911 DefOpCode = VRD->getOpcode();
1913 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1914 DefOpCode == TargetOpcode::G_CONSTANT ||
1915 DefOpCode == SPIRV::OpVariable ||
1916 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1917 DefOpCode == SPIRV::OpConstantI) {
1923 uint32_t SpecOpcode = 0;
1925 case SPIRV::OpConvertPtrToU:
1926 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1928 case SPIRV::OpConvertUToPtr:
1929 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1934 TII.get(SPIRV::OpSpecConstantOp))
1944 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1948bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1949 SPIRVTypeInst ResType,
1950 MachineInstr &
I)
const {
1951 Register OpReg =
I.getOperand(1).getReg();
1952 SPIRVTypeInst OpType =
1955 return diagnoseUnsupported(
1956 I,
"incompatible result and operand types in a bitcast");
1957 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1963 std::optional<Align> AlignOverride = std::nullopt) {
1968 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1969 if (
MemOp->isNonTemporal())
1970 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1972 if (!ST->isShader() &&
MemOp->getAlign().value())
1973 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1977 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1978 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1982 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1984 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1988 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1992 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1994 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1995 MIB.
addImm(AlignOverride.value_or(
MemOp->getAlign()).value());
2006 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2008 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2010 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2014bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2015 SPIRVTypeInst ResType,
2016 MachineInstr &
I)
const {
2018 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2023 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2024 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2026 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2028 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2032 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2036 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2037 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2038 I.getDebugLoc(),
I);
2042 MachineIRBuilder MIRBuilder(
I);
2044 if (
I.getNumMemOperands()) {
2045 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2046 if (MemOp->isAtomic())
2047 return selectAtomicLoad(ResVReg, ResType,
I);
2050 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2054 if (!
I.getNumMemOperands()) {
2055 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2057 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2066Register SPIRVInstructionSelector::createPtrSizedIntReg(
2067 MachineIRBuilder &MIRBuilder)
const {
2068 SPIRVTypeInst IntType =
2078SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2079 MachineIRBuilder &MIRBuilder)
const {
2080 SPIRVTypeInst IntType =
2082 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2083 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2091Register SPIRVInstructionSelector::castPtrToPtrToInt(
2092 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2093 MachineIRBuilder &MIRBuilder)
const {
2094 SPIRVTypeInst IntType =
2096 SPIRVTypeInst PtrType =
2110bool SPIRVInstructionSelector::selectAtomicPtrValue(
2111 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2112 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2120 Register IntResult = EmitAtomic(IntType);
2122 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2130bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2131 SPIRVTypeInst ResType,
2132 MachineInstr &
I)
const {
2133 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2136 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2139 return diagnoseUnsupported(
2140 I,
"Lowering to SPIR-V of atomic load is only "
2141 "allowed for integer, floating point or pointer types");
2143 assert(
I.getNumMemOperands());
2144 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2145 assert(MemOp.isAtomic());
2147 uint32_t
Scope =
static_cast<uint32_t
>(
2149 Register ScopeReg = buildI32Constant(Scope,
I);
2155 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2156 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2159 Register MemSemReg = buildI32Constant(Sem,
I);
2161 MachineIRBuilder MIRBuilder(
I);
2165 return diagnoseUnsupported(
2166 I,
"Lowering to SPIR-V of atomic load is only "
2167 "allowed for pointer types for physical addressing model");
2172 SPIRV::StorageClass::StorageClass SC =
2174 return selectAtomicPtrValue(
2175 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2176 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2177 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2188 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2199bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2201 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2202 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2207 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2208 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2210 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2215 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2219 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2220 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2221 SPIRVTypeInst SampledType =
2223 SPIRVTypeInst StoreValCompType =
2225 if (StoreValCompType && StoreValCompType != SampledType) {
2228 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2231 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2236 StoreVal = PackedReg;
2239 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2240 TII.get(SPIRV::OpImageWrite))
2246 if (sampledTypeIsSignedInteger(LLVMHandleType))
2249 BMI.constrainAllUses(
TII,
TRI, RBI);
2256 if (PointeeTy && PointeeTy->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
2257 StoreTy->
getOpcode() != SPIRV::OpTypeVectorIdEXT &&
2259 MachineInstr *StoreValDef =
getVRegDef(*MRI, StoreVal);
2271 if (
I.getNumMemOperands()) {
2272 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2273 if (MemOp->isAtomic())
2274 return selectAtomicStore(
I);
2281 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2282 PtrSC == SPIRV::StorageClass::Input ||
2283 PtrSC == SPIRV::StorageClass::PushConstant)
2284 return diagnoseUnsupported(
2285 I,
"store into a read-only SPIR-V storage class is not allowed");
2287 MachineIRBuilder MIRBuilder(
I);
2289 if (!
I.getNumMemOperands()) {
2290 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2292 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2301bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2302 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2305 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2306 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2311 if (!PointeeType && PtrType &&
2312 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2315 return diagnoseUnsupported(
I,
2316 "Lowering to SPIR-V of atomic store is only "
2317 "allowed for integer or floating point types");
2319 assert(
I.getNumMemOperands());
2320 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2321 assert(MemOp.isAtomic());
2323 uint32_t
Scope =
static_cast<uint32_t
>(
2325 Register ScopeReg = buildI32Constant(Scope,
I);
2331 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2332 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2335 Register MemSemReg = buildI32Constant(Sem,
I);
2336 MachineIRBuilder MIRBuilder(
I);
2340 return diagnoseUnsupported(
2341 I,
"Lowering to SPIR-V of atomic store is only "
2342 "allowed for pointer types for physical addressing model");
2347 SPIRV::StorageClass::StorageClass SC =
2349 return selectAtomicPtrValue(
2350 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2352 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2365 return diagnoseUnsupported(
I,
2366 "Lowering to SPIR-V of atomic store is only "
2367 "allowed for integer or floating point types");
2369 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2379bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2380 SPIRVTypeInst ResType,
2381 MachineInstr &
I)
const {
2382 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2390 const Register PtrsReg =
I.getOperand(2).getReg();
2391 const uint32_t
Alignment =
I.getOperand(3).getImm();
2392 const Register MaskReg =
I.getOperand(4).getReg();
2393 const Register PassthruReg =
I.getOperand(5).getReg();
2394 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2398 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2409bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2410 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2417 const Register ValuesReg =
I.getOperand(1).getReg();
2418 const Register PtrsReg =
I.getOperand(2).getReg();
2419 const uint32_t
Alignment =
I.getOperand(3).getImm();
2420 const Register MaskReg =
I.getOperand(4).getReg();
2421 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2425 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2434bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2435 const Twine &
Msg)
const {
2436 const Function &
F =
I.getMF()->getFunction();
2437 F.getContext().diagnose(
2438 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2442bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2443 SPIRVTypeInst ResType,
2444 MachineInstr &
I)
const {
2445 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2446 return diagnoseUnsupported(
2447 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2448 "SPIR-V extension: SPV_INTEL_variable_length_array");
2450 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2457bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2458 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2459 return diagnoseUnsupported(
2461 "llvm.stackrestore intrinsic: this instruction requires the following "
2462 "SPIR-V extension: SPV_INTEL_variable_length_array");
2463 if (!
I.getOperand(0).isReg())
2466 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2467 .
addUse(
I.getOperand(0).getReg())
2473SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2474 MachineIRBuilder MIRBuilder(
I);
2475 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2482 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2486 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2487 Type *ArrTy = ArrayType::get(ValTy, Num);
2489 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2492 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2503 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2504 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2505 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2506 : SPIRV::OpVariable))
2509 .
addImm(SPIRV::StorageClass::UniformConstant);
2522bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2525 Register DstReg =
I.getOperand(0).getReg();
2531 return diagnoseUnsupported(
2532 I,
"OpCopyMemory requires operands to have the same type");
2537 return diagnoseUnsupported(
2538 I,
"Unable to determine pointee type size for OpCopyMemory");
2539 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2540 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2541 return diagnoseUnsupported(
2542 I,
"OpCopyMemory requires the size to match the pointee type size");
2545 const unsigned Opcode =
2546 IsLogical ? SPIRV::OpCopyMemory : SPIRV::OpCopyMemorySized;
2547 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2551 MIB.
addUse(
I.getOperand(2).getReg());
2553 if (
I.getNumMemOperands()) {
2554 MachineIRBuilder MIRBuilder(
I);
2555 const MachineMemOperand *DstMemOp = *
I.memoperands_begin();
2557 Align SrcAlign = DstAlign;
2560 if (
I.getNumMemOperands() > 1 && !STI.
isShader())
2561 SrcAlign = (*std::next(
I.memoperands_begin()))->
getAlign();
2571 std::min(DstAlign, SrcAlign));
2578bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2579 MachineInstr &
I)
const {
2581 Register SizeReg =
I.getOperand(2).getReg();
2583 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2587 Register SrcReg =
I.getOperand(1).getReg();
2588 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2589 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2590 Register VarReg = getOrCreateMemSetGlobal(
I);
2593 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2595 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2597 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2600 if (!selectCopyMemory(
I, SrcReg))
2602 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2603 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2608bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2609 SPIRVTypeInst ResType,
2612 unsigned NegateOpcode)
const {
2614 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2615 uint32_t
Scope =
static_cast<uint32_t
>(
2617 MemOp->getSyncScopeID()));
2618 Register ScopeReg = buildI32Constant(Scope,
I);
2620 Register Ptr =
I.getOperand(1).getReg();
2621 uint32_t ScSem =
static_cast<uint32_t
>(
2625 Register MemSemReg = buildI32Constant(
2629 Register ValueReg =
I.getOperand(2).getReg();
2630 if (NegateOpcode != 0) {
2633 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2639 if (NewOpcode != SPIRV::OpAtomicExchange)
2640 return diagnoseUnsupported(
2641 I,
"Lowering to SPIR-V of this atomic operation is not "
2642 "allowed for pointer types");
2644 return diagnoseUnsupported(
2645 I,
"Lowering to SPIR-V of atomic exchange is only "
2646 "allowed for pointer types for physical addressing model");
2653 MachineIRBuilder MIRBuilder(
I);
2655 return selectAtomicPtrValue(
2656 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2658 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2659 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2660 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2668 return ExchangeResReg;
2672 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2683bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2684 unsigned ArgI =
I.getNumOperands() - 1;
2686 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2687 SPIRVTypeInst SrcType =
2691 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2695 unsigned CurrentIndex = 0;
2696 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2697 Register ResVReg =
I.getOperand(i).getReg();
2700 LLT ResLLT = MRI->
getType(ResVReg);
2706 ResType = ScalarType;
2715 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2721 for (
unsigned j = 0;
j < NumElements; ++
j) {
2722 MIB.
addImm(CurrentIndex + j);
2724 CurrentIndex += NumElements;
2728 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2740bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2743 ? SPIRV::MemorySemantics::UniformMemory |
2744 SPIRV::MemorySemantics::WorkgroupMemory |
2745 SPIRV::MemorySemantics::ImageMemory
2746 : SPIRV::MemorySemantics::WorkgroupMemory |
2747 SPIRV::MemorySemantics::CrossWorkgroupMemory |
2748 SPIRV::MemorySemantics::ImageMemory;
2750 STI.getTargetTriple(),
static_cast<uint32_t
>(
getMemSemantics(AO)), ScSem);
2751 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
2755 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2757 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2764bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2765 SPIRVTypeInst ResType,
2767 unsigned Opcode)
const {
2768 Type *ResTy =
nullptr;
2771 return diagnoseUnsupported(
2773 "Not enough info to select the arithmetic with overflow instruction");
2775 return diagnoseUnsupported(
I,
2776 "Expect struct type result for the arithmetic "
2777 "with overflow instruction");
2783 MachineIRBuilder MIRBuilder(
I);
2785 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2786 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2793 Register ZeroReg = buildZerosVal(ResType,
I);
2798 if (ResName.
size() > 0)
2806 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2807 MIB.
addUse(
I.getOperand(i).getReg());
2812 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2813 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2815 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2816 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2823 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2824 .
addDef(
I.getOperand(1).getReg())
2832bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2833 SPIRVTypeInst ResType,
2834 MachineInstr &
I)
const {
2836 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2837 Register Ptr =
I.getOperand(2).getReg();
2838 Register ScopeReg =
I.getOperand(5).getReg();
2839 Register MemSemEqReg =
I.getOperand(6).getReg();
2840 Register MemSemNeqReg =
I.getOperand(7).getReg();
2842 Register Val =
I.getOperand(4).getReg();
2846 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2865 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2872 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2884 case SPIRV::StorageClass::DeviceOnlyINTEL:
2885 case SPIRV::StorageClass::HostOnlyINTEL:
2894 bool IsGRef =
false;
2895 bool IsAllowedRefs =
2897 unsigned Opcode = It.getOpcode();
2898 if (Opcode == SPIRV::OpConstantComposite ||
2899 Opcode == SPIRV::OpSpecConstantComposite ||
2900 Opcode == SPIRV::OpVariable ||
2901 Opcode == SPIRV::OpUntypedVariableKHR ||
2902 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2903 return IsGRef = true;
2904 return Opcode == SPIRV::OpName;
2906 return IsAllowedRefs && IsGRef;
2909Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2910 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2912 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2916SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2918 uint32_t Opcode)
const {
2919 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2920 TII.get(SPIRV::OpSpecConstantOp))
2928SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2929 SPIRVTypeInst SrcPtrTy)
const {
2930 SPIRVTypeInst GenericPtrTy =
2934 SPIRV::StorageClass::Generic),
2938 MachineInstrBuilder MIB = buildSpecConstantOp(
2940 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2950bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2951 SPIRVTypeInst ResType,
2952 MachineInstr &
I)
const {
2956 Register SrcPtr =
I.getOperand(1).getReg();
2961 return BuildCOPY(ResVReg, SrcPtr,
I);
2971 unsigned SpecOpcode = [&]() ->
unsigned {
2972 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2973 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2975 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2977 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
2985 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2987 .constrainAllUses(
TII,
TRI, RBI);
2989 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2991 buildSpecConstantOp(
2993 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2994 .constrainAllUses(
TII,
TRI, RBI);
3001 return BuildCOPY(ResVReg, SrcPtr,
I);
3003 if ((SrcSC == SPIRV::StorageClass::Function &&
3004 DstSC == SPIRV::StorageClass::Private) ||
3005 (DstSC == SPIRV::StorageClass::Function &&
3006 SrcSC == SPIRV::StorageClass::Private))
3007 return BuildCOPY(ResVReg, SrcPtr,
I);
3011 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3014 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3017 SPIRVTypeInst GenericPtrTy =
3036 return selectUnOp(ResVReg, ResType,
I,
3037 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
3039 return selectUnOp(ResVReg, ResType,
I,
3040 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
3042 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3044 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3054bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3055 SPIRVTypeInst ResType,
3056 MachineInstr &
I)
const {
3058 return diagnoseUnsupported(
3059 I,
"G_PTRMASK is not supported with logical SPIR-V");
3064 Register PtrReg =
I.getOperand(1).getReg();
3065 Register MaskReg =
I.getOperand(2).getReg();
3084 ? SPIRV::OpBitwiseAndV
3085 : SPIRV::OpBitwiseAndS;
3108 return SPIRV::OpFOrdEqual;
3110 return SPIRV::OpFOrdGreaterThanEqual;
3112 return SPIRV::OpFOrdGreaterThan;
3114 return SPIRV::OpFOrdLessThanEqual;
3116 return SPIRV::OpFOrdLessThan;
3118 return SPIRV::OpFOrdNotEqual;
3120 return SPIRV::OpOrdered;
3122 return SPIRV::OpFUnordEqual;
3124 return SPIRV::OpFUnordGreaterThanEqual;
3126 return SPIRV::OpFUnordGreaterThan;
3128 return SPIRV::OpFUnordLessThanEqual;
3130 return SPIRV::OpFUnordLessThan;
3132 return SPIRV::OpFUnordNotEqual;
3134 return SPIRV::OpUnordered;
3144 return SPIRV::OpIEqual;
3146 return SPIRV::OpINotEqual;
3148 return SPIRV::OpSGreaterThanEqual;
3150 return SPIRV::OpSGreaterThan;
3152 return SPIRV::OpSLessThanEqual;
3154 return SPIRV::OpSLessThan;
3156 return SPIRV::OpUGreaterThanEqual;
3158 return SPIRV::OpUGreaterThan;
3160 return SPIRV::OpULessThanEqual;
3162 return SPIRV::OpULessThan;
3171 return SPIRV::OpPtrEqual;
3173 return SPIRV::OpPtrNotEqual;
3184 return SPIRV::OpLogicalEqual;
3186 return SPIRV::OpLogicalNotEqual;
3224bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3225 SPIRVTypeInst ResType,
3227 unsigned OpAnyOrAll)
const {
3228 assert(
I.getNumOperands() == 3);
3229 assert(
I.getOperand(2).isReg());
3231 Register InputRegister =
I.getOperand(2).getReg();
3234 assert(InputType &&
"VReg has no type assigned");
3238 assert(ResVReg ==
I.getOperand(0).getReg());
3239 return BuildCOPY(ResVReg, InputRegister,
I);
3243 unsigned SpirvNotEqualId =
3244 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3246 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3251 IsBoolTy ? InputRegister
3259 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3261 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3278bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3279 SPIRVTypeInst ResType,
3280 MachineInstr &
I)
const {
3281 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3284bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3285 SPIRVTypeInst ResType,
3286 MachineInstr &
I)
const {
3287 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3291bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3292 SPIRVTypeInst ResType,
3293 MachineInstr &
I)
const {
3294 assert(
I.getNumOperands() == 4);
3295 assert(
I.getOperand(2).isReg());
3296 assert(
I.getOperand(3).isReg());
3298 [[maybe_unused]] SPIRVTypeInst VecType =
3303 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3304 "dot product requires either a vector of at least 2 components or"
3305 " the SPV_EXT_long vector extension.");
3307 [[maybe_unused]] SPIRVTypeInst EltType =
3316 .
addUse(
I.getOperand(2).getReg())
3317 .
addUse(
I.getOperand(3).getReg())
3322bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3323 SPIRVTypeInst ResType,
3326 assert(
I.getNumOperands() == 4);
3327 assert(
I.getOperand(2).isReg());
3328 assert(
I.getOperand(3).isReg());
3331 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3335 .
addUse(
I.getOperand(2).getReg())
3336 .
addUse(
I.getOperand(3).getReg())
3343bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3344 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3345 assert(
I.getNumOperands() == 4);
3346 assert(
I.getOperand(2).isReg());
3347 assert(
I.getOperand(3).isReg());
3351 Register Vec0 =
I.getOperand(2).getReg();
3352 Register Vec1 =
I.getOperand(3).getReg();
3356 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3365 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3366 "dot product requires either a vector of at least 2 components "
3367 "or the SPV_EXT_long_vector extension.");
3370 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3380 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3391 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3403bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3404 SPIRVTypeInst ResType,
3405 MachineInstr &
I)
const {
3407 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3410 .
addUse(
I.getOperand(2).getReg())
3415bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3416 SPIRVTypeInst ResType,
3417 MachineInstr &
I)
const {
3419 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3422 .
addUse(
I.getOperand(2).getReg())
3427bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3428 SPIRVTypeInst ResType,
3429 MachineInstr &
I)
const {
3432 Register Src =
I.getOperand(2).getReg();
3475bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3476 SPIRVTypeInst ResType,
3477 MachineInstr &
I)
const {
3479 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3482 .
addUse(
I.getOperand(2).getReg())
3487template <
bool Signed>
3488bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3489 SPIRVTypeInst ResType,
3490 MachineInstr &
I)
const {
3491 assert(
I.getNumOperands() == 5);
3492 assert(
I.getOperand(2).isReg());
3493 assert(
I.getOperand(3).isReg());
3494 assert(
I.getOperand(4).isReg());
3497 Register Acc =
I.getOperand(2).getReg();
3501 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3503 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3508 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3511 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3523template <
bool Signed>
3524bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3525 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3526 assert(
I.getNumOperands() == 5);
3527 assert(
I.getOperand(2).isReg());
3528 assert(
I.getOperand(3).isReg());
3529 assert(
I.getOperand(4).isReg());
3532 Register Acc =
I.getOperand(2).getReg();
3538 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3542 for (
unsigned i = 0; i < 4; i++) {
3565 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3585 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3600bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3601 SPIRVTypeInst ResType,
3602 MachineInstr &
I)
const {
3603 assert(
I.getNumOperands() == 3);
3604 assert(
I.getOperand(2).isReg());
3606 Register VZero = buildZerosValF(ResType,
I);
3607 Register VOne = buildOnesValF(ResType,
I);
3609 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3612 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3614 .
addUse(
I.getOperand(2).getReg())
3621bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3622 SPIRVTypeInst ResType,
3623 MachineInstr &
I)
const {
3624 assert(
I.getNumOperands() == 3);
3625 assert(
I.getOperand(2).isReg());
3627 Register InputRegister =
I.getOperand(2).getReg();
3629 auto &
DL =
I.getDebugLoc();
3632 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3639 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3641 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3649 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3654 if (NeedsConversion) {
3655 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3666bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3667 SPIRVTypeInst ResType,
3669 unsigned Opcode)
const {
3673 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3679 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3680 BMI.addUse(
I.getOperand(J).getReg());
3687bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3690 bool WithGroupSync)
const {
3692 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3694 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3696 assert(((Scope != SPIRV::Scope::Workgroup) ||
3697 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3698 "Workgroup Scope must set WorkGroupMemory semantic "
3699 "in Barrier instruction");
3701 assert(((Scope != SPIRV::Scope::Device) ||
3702 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3703 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3704 "Device Scope must set UniformMemory and ImageMemory semantic "
3705 "in Barrier instruction");
3711 if (WithGroupSync) {
3712 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3716 Register ScopeReg = buildI32Constant(Scope,
I);
3717 Register MemSemReg = buildI32Constant(MemSem,
I);
3719 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3723bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3724 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3729 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3730 SPIRV::OpGroupNonUniformBallot))
3735 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3740 .
addImm(SPIRV::GroupOperation::Reduce)
3747bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3748 SPIRVTypeInst ResType,
3749 MachineInstr &
I)
const {
3754 Register InputReg =
I.getOperand(2).getReg();
3759 bool IsVector = NumElems > 1 ||
3760 (InputType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
3774 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3775 SPIRV::OpGroupNonUniformAllEqual);
3780 ElementResults.
reserve(NumElems);
3782 for (
unsigned Idx = 0;
Idx < NumElems; ++
Idx) {
3795 ElemInput = Extracted;
3801 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3812 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3823bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3824 SPIRVTypeInst ResType,
3825 MachineInstr &
I)
const {
3827 assert(
I.getNumOperands() == 3);
3829 auto Op =
I.getOperand(2);
3839 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3841 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3842 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3863 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3867 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3874bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3875 SPIRVTypeInst ResType,
3877 bool IsUnsigned)
const {
3878 return selectWaveReduce(
3879 ResVReg, ResType,
I, IsUnsigned,
3880 [&](
Register InputRegister,
bool IsUnsigned) {
3881 const bool IsFloatTy =
3883 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3884 : SPIRV::OpGroupNonUniformSMax;
3885 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3889bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3890 SPIRVTypeInst ResType,
3892 bool IsUnsigned)
const {
3893 return selectWaveReduce(
3894 ResVReg, ResType,
I, IsUnsigned,
3895 [&](
Register InputRegister,
bool IsUnsigned) {
3896 const bool IsFloatTy =
3898 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3899 : SPIRV::OpGroupNonUniformSMin;
3900 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3904bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3905 SPIRVTypeInst ResType,
3906 MachineInstr &
I)
const {
3907 return selectWaveReduce(ResVReg, ResType,
I,
false,
3908 [&](
Register InputRegister,
bool IsUnsigned) {
3910 InputRegister, SPIRV::OpTypeFloat);
3911 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3912 : SPIRV::OpGroupNonUniformIAdd;
3916bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3917 SPIRVTypeInst ResType,
3918 MachineInstr &
I)
const {
3919 return selectWaveReduce(ResVReg, ResType,
I,
false,
3920 [&](
Register InputRegister,
bool IsUnsigned) {
3922 InputRegister, SPIRV::OpTypeFloat);
3923 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3924 : SPIRV::OpGroupNonUniformIMul;
3928template <
typename PickOpcodeFn>
3929bool SPIRVInstructionSelector::selectWaveReduce(
3930 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3931 PickOpcodeFn &&PickOpcode)
const {
3932 assert(
I.getNumOperands() == 3);
3933 assert(
I.getOperand(2).isReg());
3935 Register InputRegister =
I.getOperand(2).getReg();
3939 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3942 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3948 .
addImm(SPIRV::GroupOperation::Reduce)
3949 .
addUse(
I.getOperand(2).getReg())
3954bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3955 SPIRVTypeInst ResType,
3957 unsigned Opcode)
const {
3958 return selectWaveReduce(
3959 ResVReg, ResType,
I,
false,
3960 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3963bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3964 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3965 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3966 [&](
Register InputRegister,
bool IsUnsigned) {
3968 InputRegister, SPIRV::OpTypeFloat);
3970 ? SPIRV::OpGroupNonUniformFAdd
3971 : SPIRV::OpGroupNonUniformIAdd;
3975bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3976 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3977 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3978 [&](
Register InputRegister,
bool IsUnsigned) {
3980 InputRegister, SPIRV::OpTypeFloat);
3982 ? SPIRV::OpGroupNonUniformFMul
3983 : SPIRV::OpGroupNonUniformIMul;
3987template <
typename PickOpcodeFn>
3988bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3989 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3990 PickOpcodeFn &&PickOpcode)
const {
3991 assert(
I.getNumOperands() == 3);
3992 assert(
I.getOperand(2).isReg());
3994 Register InputRegister =
I.getOperand(2).getReg();
3998 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
4001 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
4007 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
4008 .
addUse(
I.getOperand(2).getReg())
4013bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
4014 SPIRVTypeInst ResType,
4017 assert(
I.getNumOperands() == 3);
4018 assert(
I.getOperand(2).isReg());
4020 Register InputRegister =
I.getOperand(2).getReg();
4026 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
4037bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
4038 SPIRVTypeInst ResType,
4045 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
4050 : SPIRV::OpUConvert;
4052 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4055 ShiftOp = SPIRV::OpShiftRightLogicalV;
4060 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4061 TII.get(SPIRV::OpConstantComposite))
4064 for (
unsigned It = 0; It <
N; ++It)
4068 ShiftConst = CompositeReg;
4073 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
4078 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
4083 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
4088 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
4091bool SPIRVInstructionSelector::handle64BitOverflow(
4093 unsigned int Opcode,
4100 "handle64BitOverflow should only be used for integer types");
4102 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4104 MachineIRBuilder MIRBuilder(
I);
4106 SPIRVTypeInst I64x2Type =
4108 SPIRVTypeInst Vec2ResType =
4111 std::vector<Register> PartialRegs;
4113 unsigned CurrentComponent = 0;
4114 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4118 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4119 TII.get(SPIRV::OpVectorShuffle))
4124 .
addImm(CurrentComponent)
4125 .
addImm(CurrentComponent + 1);
4135 PartialRegs.push_back(SubVecReg);
4138 if (CurrentComponent != ComponentCount) {
4144 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4145 SPIRV::OpVectorExtractDynamic))
4154 PartialRegs.push_back(FinalElemResReg);
4158 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4159 SPIRV::OpCompositeConstruct);
4162bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4163 SPIRVTypeInst ResType,
4167 if (ComponentCount > 2)
4168 return handle64BitOverflow(
4169 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4171 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4173 MachineIRBuilder MIRBuilder(
I);
4177 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4181 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4186 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4193 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4194 TII.get(SPIRV::OpVectorShuffle))
4199 for (
unsigned J = 0; J < ComponentCount; ++J) {
4206 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4209bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4210 SPIRVTypeInst ResType,
4214 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4222bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4223 SPIRVTypeInst ResType,
4224 MachineInstr &
I)
const {
4225 Register OpReg =
I.getOperand(1).getReg();
4234 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4236 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4238 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4240 return SPIRVInstructionSelector::diagnoseUnsupported(
4241 I,
"G_BITREVERSE only support 16,32,64 bits.");
4245 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4256 unsigned AndOp = SPIRV::OpBitwiseAndS;
4257 unsigned OrOp = SPIRV::OpBitwiseOrS;
4258 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4259 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4260 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4262 AndOp = SPIRV::OpBitwiseAndV;
4263 OrOp = SPIRV::OpBitwiseOrV;
4264 ShlOp = SPIRV::OpShiftLeftLogicalV;
4265 ShrOp = SPIRV::OpShiftRightLogicalV;
4271 const unsigned Shift) ->
Register {
4274 (ResType->
getOpcode() != SPIRV::OpTypeVectorIdEXT ||
4281 Register MaskReg = CreateConst(Mask);
4282 Register ShiftReg = CreateConst(Shift);
4289 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4290 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4291 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4292 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4293 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4302 while ((Shift >>= 1) > 0) {
4309 return BuildCOPY(ResVReg, Result,
I);
4312bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4313 SPIRVTypeInst ResType,
4314 MachineInstr &
I)
const {
4315 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4316 "G_FREEZE must define and use a register");
4317 Register OpReg =
I.getOperand(1).getReg();
4321 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4334 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4335 if (
Def->getOpcode() == TargetOpcode::COPY)
4338 switch (
Def->getOpcode()) {
4339 case SPIRV::ASSIGN_TYPE:
4340 if (MachineInstr *AssignToDef =
4342 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4343 Reg =
Def->getOperand(2).getReg();
4346 case SPIRV::OpUndef:
4347 Reg =
Def->getOperand(1).getReg();
4350 unsigned DestOpCode;
4352 DestOpCode = SPIRV::OpConstantNull;
4353 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4354 "static undef/poison lowered to OpConstantNull\n");
4356 DestOpCode = TargetOpcode::COPY;
4358 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4359 "skipped, lowered as a copy of the operand\n");
4361 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4362 .
addDef(
I.getOperand(0).getReg())
4370bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4371 SPIRVTypeInst ResType,
4372 MachineInstr &
I)
const {
4376 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4380 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4385 for (
unsigned i =
I.getNumExplicitDefs();
4386 i <
I.getNumExplicitOperands() && IsConst; ++i)
4391 return diagnoseUnsupported(
4392 I,
"There must be at least two constituent operands in a vector");
4397 for (
unsigned i =
I.getNumExplicitDefs();
4398 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4399 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4404 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4411 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4412 TII.get(IsConst ? SPIRV::OpConstantComposite
4413 : SPIRV::OpCompositeConstruct))
4416 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4417 MIB.
addUse(
I.getOperand(i).getReg());
4422bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4423 SPIRVTypeInst ResType,
4424 MachineInstr &
I)
const {
4428 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4433 unsigned OpIdx =
I.getNumExplicitDefs();
4434 if (!
I.getOperand(OpIdx).isReg())
4438 Register OpReg =
I.getOperand(OpIdx).getReg();
4442 return diagnoseUnsupported(
4443 I,
"There must be at least two constituent operands in a vector");
4446 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4447 TII.get(IsConst ? SPIRV::OpConstantComposite
4448 : SPIRV::OpCompositeConstruct))
4451 for (
unsigned i = 0; i <
N; ++i)
4457bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4458 SPIRVTypeInst ResType,
4459 MachineInstr &
I)
const {
4465 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4467 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4468 TII.get(SPIRV::OpCompositeConstruct))
4471 for (
unsigned OpIdx =
I.getNumExplicitDefs();
4472 OpIdx <
I.getNumExplicitOperands(); ++OpIdx)
4473 MIB.
addUse(
I.getOperand(OpIdx).getReg());
4478bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4479 SPIRVTypeInst ResType,
4480 MachineInstr &
I)
const {
4486 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4488 Opcode = SPIRV::OpDemoteToHelperInvocation;
4490 Opcode = SPIRV::OpKill;
4495 ToErase.eraseFromParent();
4504bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4505 SPIRVTypeInst ResType,
unsigned CmpOpc,
4506 MachineInstr &
I)
const {
4507 Register Cmp0 =
I.getOperand(2).getReg();
4508 Register Cmp1 =
I.getOperand(3).getReg();
4511 "CMP operands should have the same type");
4512 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4522bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4523 SPIRVTypeInst ResType,
4524 MachineInstr &
I)
const {
4525 auto Pred =
I.getOperand(1).getPredicate();
4528 Register CmpOperand =
I.getOperand(2).getReg();
4530 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4535 Register Op1 =
I.getOperand(3).getReg();
4539 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4544 I.getOperand(3).setReg(NewOp1);
4550 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4554SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4555 SPIRVTypeInst ResType)
const {
4557 SPIRVTypeInst SpvI32Ty =
4560 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4567 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4570 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4573 .
addImm(APInt(32, Val).getZExtValue());
4575 GR.
add(ConstInt,
MI);
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 {
4616 auto Pred =
I.getOperand(1).getPredicate();
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 {
5107 if (
getImm(
I.getOperand(4), MRI))
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 {
5123 if (
getImm(
I.getOperand(3), MRI))
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);
5215 assert((!IsAccessChainOpcode || (
getImm(
I.getOperand(4), MRI) &&
5216 foldImm(
I.getOperand(4), MRI) == 0)) &&
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();
5263 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
5265 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
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: {
5388 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
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 {
5956 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5963bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5964 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5966 assert(Intr.getIntrinsicID() ==
5967 Intrinsic::spv_resource_counterhandlefrombinding);
5970 Register MainHandleReg = Intr.getOperand(2).getReg();
5972 assert(MainHandleDef->getIntrinsicID() ==
5973 Intrinsic::spv_resource_handlefrombinding);
5977 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5978 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5979 std::string CounterName =
5984 MachineIRBuilder MIRBuilder(
I);
5986 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5988 ArraySize, IndexReg, CounterName, MIRBuilder);
5990 return BuildCOPY(ResVReg, CounterVarReg,
I);
5993bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5994 SPIRVTypeInst ResType,
5995 MachineInstr &
I)
const {
5997 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5999 Register CounterHandleReg = Intr.getOperand(2).getReg();
6000 Register IncrReg = Intr.getOperand(3).getReg();
6007 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
6008 assert(CounterVarPointeeType &&
6009 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
6010 "Counter variable must be a struct");
6012 SPIRV::StorageClass::StorageBuffer &&
6013 "Counter variable must be in the storage buffer storage class");
6015 "Counter variable must have exactly 1 member in the struct");
6016 const SPIRVTypeInst MemberType =
6019 "Counter variable struct must have a single i32 member");
6023 MachineIRBuilder MIRBuilder(
I);
6025 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
6028 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
6034 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6037 .
addUse(CounterHandleReg)
6044 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
6047 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
6050 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
6059 return BuildCOPY(ResVReg, AtomicRes,
I);
6067 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
6075bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
6076 SPIRVTypeInst ResType,
6077 MachineInstr &
I)
const {
6085 Register ImageReg =
I.getOperand(2).getReg();
6093 Register IdxReg =
I.getOperand(3).getReg();
6095 MachineInstr &Pos =
I;
6097 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
6101bool SPIRVInstructionSelector::generateSampleImage(
6104 DebugLoc Loc, MachineInstr &Pos)
const {
6115 if (!loadHandleBeforePosition(NewSamplerReg,
6121 MachineIRBuilder MIRBuilder(Pos);
6134 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
6135 ImOps.Lod.has_value();
6136 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
6137 : SPIRV::OpImageSampleImplicitLod;
6139 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
6140 : SPIRV::OpImageSampleDrefImplicitLod;
6149 MIB.
addUse(*ImOps.Compare);
6151 uint32_t ImageOperands = 0;
6153 ImageOperands |= SPIRV::ImageOperand::Bias;
6155 ImageOperands |= SPIRV::ImageOperand::Lod;
6156 if (ImOps.GradX && ImOps.GradY)
6157 ImageOperands |= SPIRV::ImageOperand::Grad;
6158 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
6160 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6163 "Non-constant offsets are not supported in sample instructions.");
6168 ImageOperands |= SPIRV::ImageOperand::MinLod;
6170 if (ImageOperands != 0) {
6171 MIB.
addImm(ImageOperands);
6172 if (ImageOperands & SPIRV::ImageOperand::Bias)
6174 if (ImageOperands & SPIRV::ImageOperand::Lod)
6176 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6177 MIB.
addUse(*ImOps.GradX);
6178 MIB.
addUse(*ImOps.GradY);
6181 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6182 MIB.
addUse(*ImOps.Offset);
6183 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6184 MIB.
addUse(*ImOps.MinLod);
6191bool SPIRVInstructionSelector::selectImageQuerySize(
6193 std::optional<Register> LodReg)
const {
6195 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6198 "ImageReg is not an image type.");
6200 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6202 unsigned NumComponents = 0;
6204 case SPIRV::Dim::DIM_1D:
6205 case SPIRV::Dim::DIM_Buffer:
6206 NumComponents =
IsArray ? 2 : 1;
6208 case SPIRV::Dim::DIM_2D:
6209 case SPIRV::Dim::DIM_Cube:
6210 case SPIRV::Dim::DIM_Rect:
6211 NumComponents =
IsArray ? 3 : 2;
6213 case SPIRV::Dim::DIM_3D:
6217 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6222 SPIRVTypeInst ResType =
6227 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6237bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6238 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6239 Register ImageReg =
I.getOperand(2).getReg();
6246 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6249bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6250 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6251 Register ImageReg =
I.getOperand(2).getReg();
6260 Register LodReg =
I.getOperand(3).getReg();
6263 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6265 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6272 TII.get(SPIRV::OpImageQueryLevels))
6279 TII.get(SPIRV::OpCompositeConstruct))
6289bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6290 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6291 Register ImageReg =
I.getOperand(2).getReg();
6302 "OpImageQuerySamples requires a multisampled image");
6304 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6312 TII.get(SPIRV::OpImageQuerySamples))
6319 TII.get(SPIRV::OpCompositeConstruct))
6329bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6330 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6331 Register ImageReg =
I.getOperand(2).getReg();
6332 Register SamplerReg =
I.getOperand(3).getReg();
6333 Register CoordinateReg =
I.getOperand(4).getReg();
6349 if (!loadHandleBeforePosition(
6354 MachineIRBuilder MIRBuilder(
I);
6360 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6370 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6377 unsigned ExtractedIndex =
6379 Intrinsic::spv_resource_calculate_lod_unclamped
6383 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6384 TII.get(SPIRV::OpCompositeExtract))
6394bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6395 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6396 Register ImageReg =
I.getOperand(2).getReg();
6397 Register SamplerReg =
I.getOperand(3).getReg();
6398 Register CoordinateReg =
I.getOperand(4).getReg();
6399 ImageOperands ImOps;
6400 if (
I.getNumOperands() > 5)
6401 ImOps.Offset =
I.getOperand(5).getReg();
6402 if (
I.getNumOperands() > 6)
6403 ImOps.MinLod =
I.getOperand(6).getReg();
6404 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6405 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6408bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6409 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6410 Register ImageReg =
I.getOperand(2).getReg();
6411 Register SamplerReg =
I.getOperand(3).getReg();
6412 Register CoordinateReg =
I.getOperand(4).getReg();
6413 ImageOperands ImOps;
6414 ImOps.Bias =
I.getOperand(5).getReg();
6415 if (
I.getNumOperands() > 6)
6416 ImOps.Offset =
I.getOperand(6).getReg();
6417 if (
I.getNumOperands() > 7)
6418 ImOps.MinLod =
I.getOperand(7).getReg();
6419 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6420 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6423bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6424 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6425 Register ImageReg =
I.getOperand(2).getReg();
6426 Register SamplerReg =
I.getOperand(3).getReg();
6427 Register CoordinateReg =
I.getOperand(4).getReg();
6428 ImageOperands ImOps;
6429 ImOps.GradX =
I.getOperand(5).getReg();
6430 ImOps.GradY =
I.getOperand(6).getReg();
6431 if (
I.getNumOperands() > 7)
6432 ImOps.Offset =
I.getOperand(7).getReg();
6433 if (
I.getNumOperands() > 8)
6434 ImOps.MinLod =
I.getOperand(8).getReg();
6435 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6436 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6439bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6440 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6441 Register ImageReg =
I.getOperand(2).getReg();
6442 Register SamplerReg =
I.getOperand(3).getReg();
6443 Register CoordinateReg =
I.getOperand(4).getReg();
6444 ImageOperands ImOps;
6445 ImOps.Lod =
I.getOperand(5).getReg();
6446 if (
I.getNumOperands() > 6)
6447 ImOps.Offset =
I.getOperand(6).getReg();
6448 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6449 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6452bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6453 SPIRVTypeInst ResType,
6454 MachineInstr &
I)
const {
6455 Register ImageReg =
I.getOperand(2).getReg();
6456 Register SamplerReg =
I.getOperand(3).getReg();
6457 Register CoordinateReg =
I.getOperand(4).getReg();
6458 ImageOperands ImOps;
6459 ImOps.Compare =
I.getOperand(5).getReg();
6460 if (
I.getNumOperands() > 6)
6461 ImOps.Offset =
I.getOperand(6).getReg();
6462 if (
I.getNumOperands() > 7)
6463 ImOps.MinLod =
I.getOperand(7).getReg();
6464 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6465 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6468bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6469 SPIRVTypeInst ResType,
6470 MachineInstr &
I)
const {
6471 Register ImageReg =
I.getOperand(2).getReg();
6472 Register CoordinateReg =
I.getOperand(3).getReg();
6473 Register LodReg =
I.getOperand(4).getReg();
6475 ImageOperands ImOps;
6477 if (
I.getNumOperands() > 5)
6478 ImOps.Offset =
I.getOperand(5).getReg();
6490 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6491 I.getDebugLoc(),
I, &ImOps);
6494bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6495 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6496 Register ImageReg =
I.getOperand(2).getReg();
6497 Register SamplerReg =
I.getOperand(3).getReg();
6498 Register CoordinateReg =
I.getOperand(4).getReg();
6499 ImageOperands ImOps;
6500 ImOps.Compare =
I.getOperand(5).getReg();
6501 if (
I.getNumOperands() > 6)
6502 ImOps.Offset =
I.getOperand(6).getReg();
6505 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6506 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6509bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6510 SPIRVTypeInst ResType,
6511 MachineInstr &
I)
const {
6512 Register ImageReg =
I.getOperand(2).getReg();
6513 Register SamplerReg =
I.getOperand(3).getReg();
6514 Register CoordinateReg =
I.getOperand(4).getReg();
6517 "ImageReg is not an image type.");
6522 ComponentOrCompareReg =
I.getOperand(5).getReg();
6523 OffsetReg =
I.getOperand(6).getReg();
6526 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6530 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6531 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6532 Dim != SPIRV::Dim::DIM_Rect) {
6534 "Gather operations are only supported for 2D, Cube, and Rect images.");
6541 if (!loadHandleBeforePosition(
6546 MachineIRBuilder MIRBuilder(
I);
6547 SPIRVTypeInst SampledImageType =
6552 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6560 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6562 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6564 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6569 .
addUse(ComponentOrCompareReg);
6571 uint32_t ImageOperands = 0;
6572 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6573 if (Dim == SPIRV::Dim::DIM_Cube) {
6575 "Gather operations with offset are not supported for Cube images.");
6579 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6581 ImageOperands |= SPIRV::ImageOperand::Offset;
6585 if (ImageOperands != 0) {
6586 MIB.
addImm(ImageOperands);
6588 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6596bool SPIRVInstructionSelector::generateImageReadOrFetch(
6599 const ImageOperands *ImOps)
const {
6602 "ImageReg is not an image type.");
6604 bool IsSignedInteger =
6609 bool IsFetch = (SampledOp.getImm() == 1);
6611 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6612 uint32_t ImageOperandsMask = 0;
6613 if (IsSignedInteger)
6614 ImageOperandsMask |= 0x1000;
6616 if (IsFetch && ImOps) {
6618 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6619 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6621 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6623 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6627 if (ImageOperandsMask != 0) {
6628 MIB.
addImm(ImageOperandsMask);
6629 if (IsFetch && ImOps) {
6632 if (ImOps->Offset &&
6633 (ImageOperandsMask &
6634 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6635 MIB.
addUse(*ImOps->Offset);
6644 SPIRVTypeInst SampledType =
6647 SPIRVTypeInst ReadType =
6648 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6649 bool ReadTypeMatchesResult = ReadType == ResType;
6651 Register ReadReg = ReadTypeMatchesResult
6657 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6663 BMI.constrainAllUses(
TII,
TRI, RBI);
6665 if (ReadTypeMatchesResult)
6678 if (ResultSize == 1) {
6687 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6690bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6691 SPIRVTypeInst ResType,
6692 MachineInstr &
I)
const {
6693 Register ResourcePtr =
I.getOperand(2).getReg();
6695 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6704 MachineIRBuilder MIRBuilder(
I);
6709 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6715 if (
I.getNumExplicitOperands() > 3) {
6716 Register IndexReg =
I.getOperand(3).getReg();
6723bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6724 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6729bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6730 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6731 Register ObjReg =
I.getOperand(2).getReg();
6732 if (!BuildCOPY(ResVReg, ObjReg,
I))
6742 decorateUsesAsNonUniform(ResVReg);
6746void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6749 {NonUniformReg,
nullptr}};
6750 llvm::SmallSet<Register, 8> Visited;
6751 while (WorkList.
size() > 0) {
6754 if (!Visited.
insert(CurrentReg).second)
6757 bool IsDecorated =
false;
6759 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6760 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6766 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6768 if (ResultReg == CurrentReg)
6776 MachineInstr &InsertPt =
6779 SPIRV::Decoration::NonUniformEXT, {});
6784bool SPIRVInstructionSelector::extractSubvector(
6786 MachineInstr &InsertionPoint)
const {
6788 [[maybe_unused]]
uint64_t InputSize =
6791 [[maybe_unused]]
bool IsLongVectorEXT =
6793 assert((InputSize > 1 || IsLongVectorEXT) &&
"The input must be a vector.");
6794 assert((ResultSize > 1 || IsLongVectorEXT) &&
"The result must be a vector.");
6795 assert(ResultSize < InputSize &&
6796 "Cannot extract more element than there are in the input.");
6803 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6812 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6814 TII.get(SPIRV::OpCompositeConstruct))
6818 for (
Register ComponentReg : ComponentRegisters)
6819 MIB.
addUse(ComponentReg);
6824bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6825 MachineInstr &
I)
const {
6832 Register ImageReg =
I.getOperand(1).getReg();
6840 Register CoordinateReg =
I.getOperand(2).getReg();
6841 Register DataReg =
I.getOperand(3).getReg();
6844 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6852Register SPIRVInstructionSelector::buildPointerToResource(
6853 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6854 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6855 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6857 if (ArraySize == 1) {
6858 SPIRVTypeInst PtrType =
6861 "SpirvResType did not have an explicit layout.");
6866 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6867 SPIRVTypeInst VarPointerType =
6870 VarPointerType, Set,
Binding, Name, MIRBuilder);
6872 SPIRVTypeInst ResPointerType =
6885bool SPIRVInstructionSelector::selectFirstBitSet16(
6886 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6887 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6889 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6893 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6896bool SPIRVInstructionSelector::selectFirstBitSet32(
6898 unsigned BitSetOpcode)
const {
6899 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6902 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6909bool SPIRVInstructionSelector::selectFirstBitSet64(
6911 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6925 if (ComponentCount > 2) {
6926 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6928 unsigned Opcode) ->
bool {
6929 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6933 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6937 MachineIRBuilder MIRBuilder(
I);
6939 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6943 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6949 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6959 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6960 SPIRV::OpVectorExtractDynamic))
6962 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6963 SPIRV::OpVectorExtractDynamic))
6967 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6968 TII.get(SPIRV::OpVectorShuffle))
6976 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6982 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6983 TII.get(SPIRV::OpVectorShuffle))
6991 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
7011 SelectOp = SPIRV::OpSelectSISCond;
7012 AddOp = SPIRV::OpIAddS;
7020 SelectOp = SPIRV::OpSelectVIVCond;
7021 AddOp = SPIRV::OpIAddV;
7027 Register RegSecondaryOffset = Reg0;
7031 if (SwapPrimarySide) {
7032 PrimaryReg = LowReg;
7033 SecondaryReg = HighReg;
7034 RegPrimaryOffset = Reg0;
7035 RegSecondaryOffset = Reg32;
7040 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
7041 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
7046 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
7047 SPIRV::OpINotEqual))
7054 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
7055 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
7060 if (SwapPrimarySide) {
7062 if (!selectOpWithSrcs(RegAdd, ResType,
I,
7063 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
7074 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
7075 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
7080 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
7081 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
7084 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
7088bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
7089 SPIRVTypeInst ResType,
7091 bool IsSigned)
const {
7093 Register OpReg =
I.getOperand(2).getReg();
7096 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
7097 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
7101 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7103 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7105 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7108 return diagnoseUnsupported(
7110 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
7114bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
7115 SPIRVTypeInst ResType,
7116 MachineInstr &
I)
const {
7118 Register OpReg =
I.getOperand(2).getReg();
7123 unsigned ExtendOpcode = SPIRV::OpUConvert;
7124 unsigned BitSetOpcode = GL::FindILsb;
7128 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7130 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7132 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7135 return diagnoseUnsupported(
I,
7136 "spv_firstbitlow only supports 16,32,64 bits.");
7140bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
7141 SPIRVTypeInst ResType,
7142 MachineInstr &
I)
const {
7147 bool UseUntypedPointers =
7148 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7149 unsigned Opcode = UseUntypedPointers
7150 ? SPIRV::OpUntypedVariableLengthArrayINTEL
7151 : SPIRV::OpVariableLengthArrayINTEL;
7153 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
7159 if (UseUntypedPointers) {
7162 "untyped variable length array result must have a recorded element "
7167 MIB.
addUse(
I.getOperand(2).getReg());
7171 unsigned Alignment =
I.getOperand(3).getImm();
7185 while (!Worklist.
empty()) {
7187 switch (
T->getOpcode()) {
7188 case SPIRV::OpTypeInt:
7189 case SPIRV::OpTypeFloat:
7190 case SPIRV::OpTypePointer:
7192 case SPIRV::OpTypeVector:
7193 case SPIRV::OpTypeVectorIdEXT:
7194 case SPIRV::OpTypeMatrix:
7195 case SPIRV::OpTypeArray: {
7196 Register OperandReg =
T->getOperand(1).getReg();
7200 case SPIRV::OpTypeStruct:
7201 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7202 Register OperandReg =
T->getOperand(Idx).getReg();
7217 Register TypeReg = Ty->getOperand(0).getReg();
7218 if (!Visited.
insert(TypeReg).second)
7221 switch (Ty->getOpcode()) {
7222 case SPIRV::OpTypePointer:
7223 if (Ty->getOperand(1).getImm() == SPIRV::StorageClass::StorageBuffer)
7227 case SPIRV::OpTypeArray:
7228 case SPIRV::OpTypeRuntimeArray:
7231 case SPIRV::OpTypeStruct:
7232 for (
unsigned I = 1;
I < Ty->getNumOperands(); ++
I)
7248bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7249 assert(
I.getNumExplicitOperands() == 2);
7251 Register MsgReg =
I.getOperand(1).getReg();
7253 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7256 return diagnoseUnsupported(
7258 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7259 "scalar, pointer, vector, matrix, or aggregate of such types)");
7262 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7269bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7278 uint32_t MsgVal = ~0
u;
7279 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7280 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7283 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7286 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7293bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7294 SPIRVTypeInst ResType,
7295 MachineInstr &
I)
const {
7302 bool UseUntypedPointers =
7303 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7305 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7308 MachineIRBuilder MIRBuilder(
I);
7311 .
addImm(SPIRV::Extension::SPV_KHR_variable_pointers);
7313 .
addImm(SPIRV::Capability::VariablePointersStorageBuffer);
7316 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7319 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7323 if (UseUntypedPointers) {
7327 return diagnoseUnsupported(
7328 I,
"could not deduce the data type of an untyped variable");
7334 unsigned Alignment =
I.getOperand(2).getImm();
7341bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7346 const MachineInstr *PrevI =
I.getPrevNode();
7348 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7352 .
addMBB(
I.getOperand(0).getMBB())
7357 .
addMBB(
I.getOperand(0).getMBB())
7362bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7373 const MachineInstr *NextI =
I.getNextNode();
7375 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7381 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7383 .
addUse(
I.getOperand(0).getReg())
7384 .
addMBB(
I.getOperand(1).getMBB())
7390bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7391 MachineInstr &
I)
const {
7393 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7395 const unsigned NumOps =
I.getNumOperands();
7396 for (
unsigned i = 1; i <
NumOps; i += 2) {
7397 MIB.
addUse(
I.getOperand(i + 0).getReg());
7398 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7404bool SPIRVInstructionSelector::selectGlobalValue(
7405 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7407 MachineIRBuilder MIRBuilder(
I);
7408 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7411 std::string GlobalIdent;
7413 unsigned &
ID = UnnamedGlobalIDs[GV];
7415 ID = UnnamedGlobalIDs.
size();
7416 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7442 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7449 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7454 MachineInstrBuilder MIB1 =
7455 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7458 MachineInstrBuilder MIB2 =
7460 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7464 GR.
add(ConstVal, MIB2);
7472 MachineInstrBuilder MIB3 =
7473 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7476 GR.
add(ConstVal, MIB3);
7482 assert(NewReg != ResVReg);
7483 return BuildCOPY(ResVReg, NewReg,
I);
7493 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7496 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7502 SPIRVTypeInst ResType =
7506 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7511 if (
GlobalVar->isExternallyInitialized() &&
7512 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7513 constexpr unsigned ReadWriteINTEL = 3u;
7516 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7522bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7523 SPIRVTypeInst ResType,
7524 MachineInstr &
I)
const {
7526 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7534 MachineIRBuilder MIRBuilder(
I);
7539 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7542 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7544 .
add(
I.getOperand(1))
7558 APFloat::rmNearestTiesToEven, &LosesInfo);
7563 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
7573bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7574 SPIRVTypeInst ResType,
7575 MachineInstr &
I)
const {
7578 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7584 Register ExpReg =
I.getOperand(2).getReg();
7586 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7587 SPIRV::OpConvertSToF))
7589 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7596bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7597 SPIRVTypeInst ResType,
7598 MachineInstr &
I)
const {
7614 MachineIRBuilder MIRBuilder(
I);
7615 SPIRVTypeInst FloatType =
7619 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7632 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7633 const bool IsUntyped =
7634 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7636 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7637 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7638 : SPIRV::OpVariable))
7641 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7649 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7652 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7655 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7659 Register IntegralPartReg =
I.getOperand(1).getReg();
7662 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7672 assert(
false &&
"GLSL::Modf is deprecated.");
7683bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7684 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7685 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7686 MachineIRBuilder MIRBuilder(
I);
7687 const SPIRVTypeInst Vec3Ty =
7690 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7702 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7706 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7712 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7719 assert(
I.getOperand(2).isReg());
7720 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7724 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7735bool SPIRVInstructionSelector::loadBuiltinInputID(
7736 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7737 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7738 MachineIRBuilder MIRBuilder(
I);
7740 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7755 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7759 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7768SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7769 MachineInstr &
I)
const {
7770 MachineIRBuilder MIRBuilder(
I);
7781bool SPIRVInstructionSelector::loadHandleBeforePosition(
7782 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7783 MachineInstr &Pos)
const {
7786 Intrinsic::spv_resource_handlefrombinding);
7794 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7795 MachineIRBuilder MIRBuilder(HandleDef);
7796 SPIRVTypeInst VarType = ResType;
7797 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7799 if (IsStructuredBuffer) {
7808 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7811 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7812 ArraySize, IndexReg, Name, MIRBuilder);
7816 uint32_t LoadOpcode =
7817 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7827bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7828 MachineInstr &
I)
const {
7830 return diagnoseUnsupported(
7831 I,
"this instruction is only supported in shaders.");
7836InstructionSelector *
7840 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
MachineInstrBuilder & UseMI
MachineInstrBuilder MachineInstrBuilder & DefMI
#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)
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)
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.
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)
int64_t getIConstValSext(Register ConstReg, const MachineRegisterInfo *MRI)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
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