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 case TargetOpcode::G_FMINIMUM:
1179 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1180 case TargetOpcode::G_FMAXNUM:
1181 case TargetOpcode::G_FMAXIMUM:
1182 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1184 case TargetOpcode::G_FCOPYSIGN:
1185 return selectCopySign(ResVReg, ResType,
I);
1187 case TargetOpcode::G_FCEIL:
1188 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1189 case TargetOpcode::G_FFLOOR:
1190 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1192 case TargetOpcode::G_FCOS:
1193 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1194 case TargetOpcode::G_FSIN:
1195 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1196 case TargetOpcode::G_FTAN:
1197 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1198 case TargetOpcode::G_FACOS:
1199 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1200 case TargetOpcode::G_FASIN:
1201 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1202 case TargetOpcode::G_FATAN:
1203 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1204 case TargetOpcode::G_FATAN2:
1205 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1206 case TargetOpcode::G_FCOSH:
1207 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1208 case TargetOpcode::G_FSINH:
1209 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1210 case TargetOpcode::G_FTANH:
1211 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1213 case TargetOpcode::G_STRICT_FSQRT:
1214 case TargetOpcode::G_FSQRT:
1215 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1217 case TargetOpcode::G_CTTZ:
1218 case TargetOpcode::G_CTTZ_ZERO_POISON:
1219 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1220 case TargetOpcode::G_CTLZ:
1221 case TargetOpcode::G_CTLZ_ZERO_POISON:
1222 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1224 case TargetOpcode::G_INTRINSIC_ROUND:
1225 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1226 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1227 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1228 case TargetOpcode::G_INTRINSIC_TRUNC:
1229 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1230 case TargetOpcode::G_FRINT:
1231 case TargetOpcode::G_FNEARBYINT:
1232 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1234 case TargetOpcode::G_SMULH:
1235 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1236 case TargetOpcode::G_UMULH:
1237 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1239 case TargetOpcode::G_SADDSAT:
1240 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1241 case TargetOpcode::G_UADDSAT:
1242 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1243 case TargetOpcode::G_SSUBSAT:
1244 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1245 case TargetOpcode::G_USUBSAT:
1246 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1248 case TargetOpcode::G_FFREXP:
1249 return selectFrexp(ResVReg, ResType,
I);
1251 case TargetOpcode::G_UADDO:
1252 return selectOverflowArith(ResVReg, ResType,
I,
1254 : SPIRV::OpIAddCarryS);
1255 case TargetOpcode::G_USUBO:
1256 return selectOverflowArith(ResVReg, ResType,
I,
1258 : SPIRV::OpISubBorrowS);
1259 case TargetOpcode::G_UMULO:
1260 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1261 case TargetOpcode::G_SMULO:
1262 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1264 case TargetOpcode::G_SEXT:
1265 return selectExt(ResVReg, ResType,
I,
true);
1266 case TargetOpcode::G_ANYEXT:
1267 case TargetOpcode::G_ZEXT:
1268 return selectExt(ResVReg, ResType,
I,
false);
1269 case TargetOpcode::G_TRUNC:
1270 return selectTrunc(ResVReg, ResType,
I);
1271 case TargetOpcode::G_FPTRUNC:
1272 case TargetOpcode::G_FPEXT:
1273 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1275 case TargetOpcode::G_PTRTOINT:
1276 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1277 case TargetOpcode::G_INTTOPTR:
1278 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1279 case TargetOpcode::G_BITCAST:
1280 return selectBitcast(ResVReg, ResType,
I);
1281 case TargetOpcode::G_ADDRSPACE_CAST:
1282 return selectAddrSpaceCast(ResVReg, ResType,
I);
1283 case TargetOpcode::G_PTRMASK:
1284 return selectPtrMask(ResVReg, ResType,
I);
1285 case TargetOpcode::G_PTR_ADD: {
1287 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1291 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1292 (*II).getOpcode() == TargetOpcode::COPY ||
1293 (*II).getOpcode() == SPIRV::OpVariable ||
1294 (*II).getOpcode() == SPIRV::OpUntypedVariableKHR) &&
1295 getImm(
I.getOperand(2), MRI));
1297 bool IsGVInit =
false;
1301 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1302 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1303 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1304 (*UseIt).getOpcode() == SPIRV::OpVariable ||
1305 (*UseIt).getOpcode() == SPIRV::OpUntypedVariableKHR) {
1317 const bool UseUntypedPointers =
1318 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1319 if (UseUntypedPointers) {
1320 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1323 .
addImm(
static_cast<uint32_t
>(
1324 SPIRV::Opcode::UntypedInBoundsPtrAccessChainKHR))
1327 .
addUse(
I.getOperand(2).getReg())
1334 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1346 return diagnoseUnsupported(
1347 I,
"incompatible result and operand types in a bitcast");
1349 MachineInstrBuilder MIB =
1350 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1357 : SPIRV::OpInBoundsPtrAccessChain))
1361 .
addUse(
I.getOperand(2).getReg())
1364 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1368 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1370 .
addUse(
I.getOperand(2).getReg())
1379 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1382 .
addImm(
static_cast<uint32_t
>(
1383 SPIRV::Opcode::InBoundsPtrAccessChain))
1386 .
addUse(
I.getOperand(2).getReg());
1391 case TargetOpcode::G_ATOMICRMW_OR:
1392 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1393 case TargetOpcode::G_ATOMICRMW_ADD:
1394 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1395 case TargetOpcode::G_ATOMICRMW_AND:
1396 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1397 case TargetOpcode::G_ATOMICRMW_MAX:
1398 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1399 case TargetOpcode::G_ATOMICRMW_MIN:
1400 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1401 case TargetOpcode::G_ATOMICRMW_SUB:
1402 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1403 case TargetOpcode::G_ATOMICRMW_XOR:
1404 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1405 case TargetOpcode::G_ATOMICRMW_UMAX:
1406 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1407 case TargetOpcode::G_ATOMICRMW_UMIN:
1408 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1409 case TargetOpcode::G_ATOMICRMW_XCHG:
1410 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1412 case TargetOpcode::G_ATOMICRMW_FADD:
1413 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1414 case TargetOpcode::G_ATOMICRMW_FSUB:
1416 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1418 : SPIRV::OpFNegate);
1419 case TargetOpcode::G_ATOMICRMW_FMIN:
1420 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1421 case TargetOpcode::G_ATOMICRMW_FMAX:
1422 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1424 case TargetOpcode::G_FENCE:
1425 return selectFence(
I);
1427 case TargetOpcode::G_STACKSAVE:
1428 return selectStackSave(ResVReg, ResType,
I);
1429 case TargetOpcode::G_STACKRESTORE:
1430 return selectStackRestore(
I);
1432 case TargetOpcode::G_UNMERGE_VALUES:
1435 case TargetOpcode::G_TRAP:
1436 case TargetOpcode::G_UBSANTRAP:
1437 return selectTrap(
I);
1442 case TargetOpcode::DBG_LABEL:
1444 case TargetOpcode::G_DEBUGTRAP:
1445 return selectDebugTrap(ResVReg, ResType,
I);
1446 case TargetOpcode::G_PREFETCH:
1447 return selectPrefetch(
I);
1454bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1455 SPIRVTypeInst ResType,
1456 MachineInstr &
I)
const {
1457 unsigned Opcode = SPIRV::OpNop;
1464bool SPIRVInstructionSelector::selectPrefetch(MachineInstr &
I)
const {
1473 MachineIRBuilder MIRBuilder(
I);
1475 const SPIRVTypeInst PointerSizeType =
1483 Register AddrVal =
I.getOperand(0).getReg();
1486 return selectExtInst(ExtReg, GR.
getOpTypeVoid(MIRBuilder),
I, CL::prefetch,
1488 {AddrVal, ConstIntOne});
1493bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1494 SPIRVTypeInst ResType,
1496 GL::GLSLExtInst GLInst,
1497 bool setMIFlags,
bool useMISrc,
1500 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1501 return diagnoseUnsupported(
1503 "this instruction is only supported with the GLSL extended instruction "
1505 return selectExtInst(ResVReg, ResType,
I,
1506 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1507 setMIFlags, useMISrc, SrcRegs);
1510bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1511 SPIRVTypeInst ResType,
1513 CL::OpenCLExtInst CLInst,
1514 bool setMIFlags,
bool useMISrc,
1516 return selectExtInst(ResVReg, ResType,
I,
1517 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1518 setMIFlags, useMISrc, SrcRegs);
1521bool SPIRVInstructionSelector::selectExtInst(
1522 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1523 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1525 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1526 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1527 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1531bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1532 SPIRVTypeInst ResType,
1535 bool setMIFlags,
bool useMISrc,
1538 for (
const auto &[InstructionSet, Opcode] : Insts) {
1542 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1545 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1550 const unsigned NumOps =
I.getNumOperands();
1553 I.getOperand(Index).getType() ==
1554 MachineOperand::MachineOperandType::MO_IntrinsicID)
1557 MIB.
add(
I.getOperand(Index));
1569bool SPIRVInstructionSelector::selectCopySign(
Register ResVReg,
1570 SPIRVTypeInst ResType,
1571 MachineInstr &
I)
const {
1573 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1578 Register MagnitudeReg =
I.getOperand(1).getReg();
1579 Register SignReg =
I.getOperand(2).getReg();
1587 unsigned AndOpcode, OrOpcode;
1588 if (ComponentCount > 1) {
1592 AndOpcode = SPIRV::OpBitwiseAndV;
1593 OrOpcode = SPIRV::OpBitwiseOrV;
1597 AndOpcode = SPIRV::OpBitwiseAndS;
1598 OrOpcode = SPIRV::OpBitwiseOrS;
1604 return selectOpWithSrcs(ResReg, IntType,
I, SrcRegs, Opcode);
1607 Register MagnitudeInt, SignInt, MagnitudeBits, SignBits, CombinedInt;
1608 if (!EmitBitOp(MagnitudeInt, {MagnitudeReg}, SPIRV::OpBitcast) ||
1609 !EmitBitOp(SignInt, {SignReg}, SPIRV::OpBitcast) ||
1610 !EmitBitOp(MagnitudeBits, {MagnitudeInt, NotSignMask}, AndOpcode) ||
1611 !EmitBitOp(SignBits, {SignInt, SignMask}, AndOpcode) ||
1612 !EmitBitOp(CombinedInt, {MagnitudeBits, SignBits}, OrOpcode))
1615 return selectOpWithSrcs(ResVReg, ResType,
I, {CombinedInt}, SPIRV::OpBitcast);
1618bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1619 SPIRVTypeInst ResType,
1620 MachineInstr &
I)
const {
1621 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1622 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1623 for (
const auto &Ex : ExtInsts) {
1624 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1625 uint32_t Opcode = Ex.second;
1629 MachineIRBuilder MIRBuilder(
I);
1632 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1639 const bool IsUntyped =
1640 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1642 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1643 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1644 : SPIRV::OpVariable))
1647 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1653 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1656 .
addImm(
static_cast<uint32_t
>(Ex.first))
1658 .
add(
I.getOperand(2))
1662 Register ExpResReg =
I.getOperand(1).getReg();
1664 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1674bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1675 SPIRVTypeInst ResType,
1676 MachineInstr &
I)
const {
1677 Register XReg =
I.getOperand(1).getReg();
1678 Register ExpReg =
I.getOperand(2).getReg();
1684 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1685 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1687 ExpReg = buildVectorSplat(ExpReg, NumElts,
I);
1690 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1691 true,
false, {XReg, ExpReg});
1694bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1695 SPIRVTypeInst ResType,
1696 MachineInstr &
I)
const {
1697 Register CosResVReg =
I.getOperand(1).getReg();
1698 unsigned SrcIdx =
I.getNumExplicitDefs();
1703 MachineIRBuilder MIRBuilder(
I);
1705 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1712 const bool IsUntyped =
1713 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1715 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1716 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1717 : SPIRV::OpVariable))
1720 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1724 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1727 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1729 .
add(
I.getOperand(SrcIdx))
1733 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1741 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1744 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1746 .
add(
I.getOperand(SrcIdx))
1748 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1751 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1753 .
add(
I.getOperand(SrcIdx))
1760bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1761 SPIRVTypeInst ResType,
1764 unsigned Opcode)
const {
1765 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1775bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1776 SPIRVTypeInst ResType,
1779 unsigned Opcode)
const {
1780 MachineIRBuilder MIRBuilder(
I);
1782 Register OpReg =
I.getOperand(1).getReg();
1788 SPIRVTypeInst ExtType =
1795 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1799 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1802 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1805bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1806 SPIRVTypeInst ResType,
1809 unsigned Opcode)
const {
1810 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1813bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1814 SPIRVTypeInst ResType,
1817 unsigned Opcode)
const {
1818 MachineIRBuilder MIRBuilder(
I);
1827 SPIRVTypeInst WorkingType =
1834 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {SrcReg}, SPIRV::OpUConvert))
1838 if (!selectOpWithSrcs(LowCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1846 IsScalar ? SPIRV::OpShiftRightLogicalS : SPIRV::OpShiftRightLogicalV;
1848 if (!selectOpWithSrcs(Shift, SrcType,
I, {SrcReg, ShiftAmount}, ShiftOp))
1852 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {Shift}, SPIRV::OpUConvert))
1856 if (!selectOpWithSrcs(HighCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1861 if (!selectOpWithSrcs(Sum, WorkingType,
I, {HighCount, LowCount},
1862 IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV))
1866 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1867 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1870bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1871 SPIRVTypeInst ResType,
1873 unsigned Opcode)
const {
1878 if (!STI.getTargetTriple().isVulkanOS())
1879 return selectUnOp(ResVReg, ResType,
I, Opcode);
1881 Register OpReg =
I.getOperand(1).getReg();
1884 : SPIRV::OpUConvert;
1888 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1890 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1892 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1894 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1898bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1899 SPIRVTypeInst ResType,
1901 unsigned Opcode)
const {
1903 Register SrcReg =
I.getOperand(1).getReg();
1908 unsigned DefOpCode = DefIt->getOpcode();
1909 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1912 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1913 DefOpCode = VRD->getOpcode();
1915 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1916 DefOpCode == TargetOpcode::G_CONSTANT ||
1917 DefOpCode == SPIRV::OpVariable ||
1918 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1919 DefOpCode == SPIRV::OpConstantI) {
1925 uint32_t SpecOpcode = 0;
1927 case SPIRV::OpConvertPtrToU:
1928 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1930 case SPIRV::OpConvertUToPtr:
1931 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1936 TII.get(SPIRV::OpSpecConstantOp))
1946 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1950bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1951 SPIRVTypeInst ResType,
1952 MachineInstr &
I)
const {
1953 Register OpReg =
I.getOperand(1).getReg();
1954 SPIRVTypeInst OpType =
1957 return diagnoseUnsupported(
1958 I,
"incompatible result and operand types in a bitcast");
1959 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1965 std::optional<Align> AlignOverride = std::nullopt) {
1970 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1971 if (
MemOp->isNonTemporal())
1972 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1974 if (!ST->isShader() &&
MemOp->getAlign().value())
1975 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1979 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1980 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1984 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1986 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1990 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1994 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1996 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1997 MIB.
addImm(AlignOverride.value_or(
MemOp->getAlign()).value());
2008 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2010 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2012 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2016bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2017 SPIRVTypeInst ResType,
2018 MachineInstr &
I)
const {
2020 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2025 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2026 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2028 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2030 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2034 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2038 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2039 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2040 I.getDebugLoc(),
I);
2044 MachineIRBuilder MIRBuilder(
I);
2046 if (
I.getNumMemOperands()) {
2047 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2048 if (MemOp->isAtomic())
2049 return selectAtomicLoad(ResVReg, ResType,
I);
2052 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2056 if (!
I.getNumMemOperands()) {
2057 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2059 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2068Register SPIRVInstructionSelector::createPtrSizedIntReg(
2069 MachineIRBuilder &MIRBuilder)
const {
2070 SPIRVTypeInst IntType =
2080SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2081 MachineIRBuilder &MIRBuilder)
const {
2082 SPIRVTypeInst IntType =
2084 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2085 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2093Register SPIRVInstructionSelector::castPtrToPtrToInt(
2094 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2095 MachineIRBuilder &MIRBuilder)
const {
2096 SPIRVTypeInst IntType =
2098 SPIRVTypeInst PtrType =
2112bool SPIRVInstructionSelector::selectAtomicPtrValue(
2113 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2114 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2122 Register IntResult = EmitAtomic(IntType);
2124 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2132bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2133 SPIRVTypeInst ResType,
2134 MachineInstr &
I)
const {
2135 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2138 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2141 return diagnoseUnsupported(
2142 I,
"Lowering to SPIR-V of atomic load is only "
2143 "allowed for integer, floating point or pointer types");
2145 assert(
I.getNumMemOperands());
2146 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2147 assert(MemOp.isAtomic());
2149 uint32_t
Scope =
static_cast<uint32_t
>(
2151 Register ScopeReg = buildI32Constant(Scope,
I);
2157 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2158 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2161 Register MemSemReg = buildI32Constant(Sem,
I);
2163 MachineIRBuilder MIRBuilder(
I);
2167 return diagnoseUnsupported(
2168 I,
"Lowering to SPIR-V of atomic load is only "
2169 "allowed for pointer types for physical addressing model");
2174 SPIRV::StorageClass::StorageClass SC =
2176 return selectAtomicPtrValue(
2177 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2178 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2179 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2190 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2201bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2203 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2204 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2209 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2210 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2212 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2217 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2221 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2222 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2223 SPIRVTypeInst SampledType =
2225 SPIRVTypeInst StoreValCompType =
2227 if (StoreValCompType && StoreValCompType != SampledType) {
2230 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2233 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2238 StoreVal = PackedReg;
2241 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2242 TII.get(SPIRV::OpImageWrite))
2248 if (sampledTypeIsSignedInteger(LLVMHandleType))
2251 BMI.constrainAllUses(
TII,
TRI, RBI);
2258 if (PointeeTy && PointeeTy->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
2259 StoreTy->
getOpcode() != SPIRV::OpTypeVectorIdEXT &&
2261 MachineInstr *StoreValDef =
getVRegDef(*MRI, StoreVal);
2273 if (
I.getNumMemOperands()) {
2274 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2275 if (MemOp->isAtomic())
2276 return selectAtomicStore(
I);
2283 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2284 PtrSC == SPIRV::StorageClass::Input ||
2285 PtrSC == SPIRV::StorageClass::PushConstant)
2286 return diagnoseUnsupported(
2287 I,
"store into a read-only SPIR-V storage class is not allowed");
2289 MachineIRBuilder MIRBuilder(
I);
2291 if (!
I.getNumMemOperands()) {
2292 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2294 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2303bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2304 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2307 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2308 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2313 if (!PointeeType && PtrType &&
2314 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2317 return diagnoseUnsupported(
I,
2318 "Lowering to SPIR-V of atomic store is only "
2319 "allowed for integer or floating point types");
2321 assert(
I.getNumMemOperands());
2322 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2323 assert(MemOp.isAtomic());
2325 uint32_t
Scope =
static_cast<uint32_t
>(
2327 Register ScopeReg = buildI32Constant(Scope,
I);
2333 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2334 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2337 Register MemSemReg = buildI32Constant(Sem,
I);
2338 MachineIRBuilder MIRBuilder(
I);
2342 return diagnoseUnsupported(
2343 I,
"Lowering to SPIR-V of atomic store is only "
2344 "allowed for pointer types for physical addressing model");
2349 SPIRV::StorageClass::StorageClass SC =
2351 return selectAtomicPtrValue(
2352 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2354 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2367 return diagnoseUnsupported(
I,
2368 "Lowering to SPIR-V of atomic store is only "
2369 "allowed for integer or floating point types");
2371 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2381bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2382 SPIRVTypeInst ResType,
2383 MachineInstr &
I)
const {
2384 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2392 const Register PtrsReg =
I.getOperand(2).getReg();
2393 const uint32_t
Alignment =
I.getOperand(3).getImm();
2394 const Register MaskReg =
I.getOperand(4).getReg();
2395 const Register PassthruReg =
I.getOperand(5).getReg();
2396 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2400 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2411bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2412 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2419 const Register ValuesReg =
I.getOperand(1).getReg();
2420 const Register PtrsReg =
I.getOperand(2).getReg();
2421 const uint32_t
Alignment =
I.getOperand(3).getImm();
2422 const Register MaskReg =
I.getOperand(4).getReg();
2423 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2427 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2436bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2437 const Twine &
Msg)
const {
2438 const Function &
F =
I.getMF()->getFunction();
2439 F.getContext().diagnose(
2440 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2444bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2445 SPIRVTypeInst ResType,
2446 MachineInstr &
I)
const {
2447 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2448 return diagnoseUnsupported(
2449 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2450 "SPIR-V extension: SPV_INTEL_variable_length_array");
2452 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2459bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2460 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2461 return diagnoseUnsupported(
2463 "llvm.stackrestore intrinsic: this instruction requires the following "
2464 "SPIR-V extension: SPV_INTEL_variable_length_array");
2465 if (!
I.getOperand(0).isReg())
2468 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2469 .
addUse(
I.getOperand(0).getReg())
2475SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2476 MachineIRBuilder MIRBuilder(
I);
2477 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2484 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2488 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2489 Type *ArrTy = ArrayType::get(ValTy, Num);
2491 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2494 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2505 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2506 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2507 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2508 : SPIRV::OpVariable))
2511 .
addImm(SPIRV::StorageClass::UniformConstant);
2524bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2527 Register DstReg =
I.getOperand(0).getReg();
2533 return diagnoseUnsupported(
2534 I,
"OpCopyMemory requires operands to have the same type");
2539 return diagnoseUnsupported(
2540 I,
"Unable to determine pointee type size for OpCopyMemory");
2541 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2542 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2543 return diagnoseUnsupported(
2544 I,
"OpCopyMemory requires the size to match the pointee type size");
2547 const unsigned Opcode =
2548 IsLogical ? SPIRV::OpCopyMemory : SPIRV::OpCopyMemorySized;
2549 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2553 MIB.
addUse(
I.getOperand(2).getReg());
2555 if (
I.getNumMemOperands()) {
2556 MachineIRBuilder MIRBuilder(
I);
2557 const MachineMemOperand *DstMemOp = *
I.memoperands_begin();
2559 Align SrcAlign = DstAlign;
2562 if (
I.getNumMemOperands() > 1 && !STI.
isShader())
2563 SrcAlign = (*std::next(
I.memoperands_begin()))->
getAlign();
2573 std::min(DstAlign, SrcAlign));
2580bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2581 MachineInstr &
I)
const {
2583 Register SizeReg =
I.getOperand(2).getReg();
2585 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2589 Register SrcReg =
I.getOperand(1).getReg();
2590 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2591 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2592 Register VarReg = getOrCreateMemSetGlobal(
I);
2595 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2597 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2599 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2602 if (!selectCopyMemory(
I, SrcReg))
2604 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2605 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2610bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2611 SPIRVTypeInst ResType,
2614 unsigned NegateOpcode)
const {
2616 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2617 uint32_t
Scope =
static_cast<uint32_t
>(
2619 MemOp->getSyncScopeID()));
2620 Register ScopeReg = buildI32Constant(Scope,
I);
2622 Register Ptr =
I.getOperand(1).getReg();
2623 uint32_t ScSem =
static_cast<uint32_t
>(
2627 Register MemSemReg = buildI32Constant(
2631 Register ValueReg =
I.getOperand(2).getReg();
2632 if (NegateOpcode != 0) {
2635 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2641 if (NewOpcode != SPIRV::OpAtomicExchange)
2642 return diagnoseUnsupported(
2643 I,
"Lowering to SPIR-V of this atomic operation is not "
2644 "allowed for pointer types");
2646 return diagnoseUnsupported(
2647 I,
"Lowering to SPIR-V of atomic exchange is only "
2648 "allowed for pointer types for physical addressing model");
2655 MachineIRBuilder MIRBuilder(
I);
2657 return selectAtomicPtrValue(
2658 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2660 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2661 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2662 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2670 return ExchangeResReg;
2674 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2685bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2686 unsigned ArgI =
I.getNumOperands() - 1;
2688 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2689 SPIRVTypeInst SrcType =
2693 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2697 unsigned CurrentIndex = 0;
2698 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2699 Register ResVReg =
I.getOperand(i).getReg();
2702 LLT ResLLT = MRI->
getType(ResVReg);
2708 ResType = ScalarType;
2717 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2723 for (
unsigned j = 0;
j < NumElements; ++
j) {
2724 MIB.
addImm(CurrentIndex + j);
2726 CurrentIndex += NumElements;
2730 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2742bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2745 ? SPIRV::MemorySemantics::UniformMemory |
2746 SPIRV::MemorySemantics::WorkgroupMemory |
2747 SPIRV::MemorySemantics::ImageMemory
2748 : SPIRV::MemorySemantics::WorkgroupMemory |
2749 SPIRV::MemorySemantics::CrossWorkgroupMemory |
2750 SPIRV::MemorySemantics::ImageMemory;
2752 STI.getTargetTriple(),
static_cast<uint32_t
>(
getMemSemantics(AO)), ScSem);
2753 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
2757 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2759 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2766bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2767 SPIRVTypeInst ResType,
2769 unsigned Opcode)
const {
2770 Type *ResTy =
nullptr;
2773 return diagnoseUnsupported(
2775 "Not enough info to select the arithmetic with overflow instruction");
2777 return diagnoseUnsupported(
I,
2778 "Expect struct type result for the arithmetic "
2779 "with overflow instruction");
2785 MachineIRBuilder MIRBuilder(
I);
2787 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2788 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2795 Register ZeroReg = buildZerosVal(ResType,
I);
2800 if (ResName.
size() > 0)
2808 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2809 MIB.
addUse(
I.getOperand(i).getReg());
2814 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2815 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2817 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2818 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2825 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2826 .
addDef(
I.getOperand(1).getReg())
2834bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2835 SPIRVTypeInst ResType,
2836 MachineInstr &
I)
const {
2838 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2839 Register Ptr =
I.getOperand(2).getReg();
2840 Register ScopeReg =
I.getOperand(5).getReg();
2841 Register MemSemEqReg =
I.getOperand(6).getReg();
2842 Register MemSemNeqReg =
I.getOperand(7).getReg();
2844 Register Val =
I.getOperand(4).getReg();
2848 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2867 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2874 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2886 case SPIRV::StorageClass::DeviceOnlyINTEL:
2887 case SPIRV::StorageClass::HostOnlyINTEL:
2896 bool IsGRef =
false;
2897 bool IsAllowedRefs =
2899 unsigned Opcode = It.getOpcode();
2900 if (Opcode == SPIRV::OpConstantComposite ||
2901 Opcode == SPIRV::OpSpecConstantComposite ||
2902 Opcode == SPIRV::OpVariable ||
2903 Opcode == SPIRV::OpUntypedVariableKHR ||
2904 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2905 return IsGRef = true;
2906 return Opcode == SPIRV::OpName;
2908 return IsAllowedRefs && IsGRef;
2911Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2912 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2914 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2918SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2920 uint32_t Opcode)
const {
2921 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2922 TII.get(SPIRV::OpSpecConstantOp))
2930SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2931 SPIRVTypeInst SrcPtrTy)
const {
2932 SPIRVTypeInst GenericPtrTy =
2936 SPIRV::StorageClass::Generic),
2940 MachineInstrBuilder MIB = buildSpecConstantOp(
2942 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2952bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2953 SPIRVTypeInst ResType,
2954 MachineInstr &
I)
const {
2958 Register SrcPtr =
I.getOperand(1).getReg();
2963 return BuildCOPY(ResVReg, SrcPtr,
I);
2973 unsigned SpecOpcode = [&]() ->
unsigned {
2974 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2975 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2977 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2979 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
2987 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2989 .constrainAllUses(
TII,
TRI, RBI);
2991 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2993 buildSpecConstantOp(
2995 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2996 .constrainAllUses(
TII,
TRI, RBI);
3003 return BuildCOPY(ResVReg, SrcPtr,
I);
3005 if ((SrcSC == SPIRV::StorageClass::Function &&
3006 DstSC == SPIRV::StorageClass::Private) ||
3007 (DstSC == SPIRV::StorageClass::Function &&
3008 SrcSC == SPIRV::StorageClass::Private))
3009 return BuildCOPY(ResVReg, SrcPtr,
I);
3013 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3016 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3019 SPIRVTypeInst GenericPtrTy =
3038 return selectUnOp(ResVReg, ResType,
I,
3039 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
3041 return selectUnOp(ResVReg, ResType,
I,
3042 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
3044 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3046 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3056bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3057 SPIRVTypeInst ResType,
3058 MachineInstr &
I)
const {
3060 return diagnoseUnsupported(
3061 I,
"G_PTRMASK is not supported with logical SPIR-V");
3066 Register PtrReg =
I.getOperand(1).getReg();
3067 Register MaskReg =
I.getOperand(2).getReg();
3086 ? SPIRV::OpBitwiseAndV
3087 : SPIRV::OpBitwiseAndS;
3110 return SPIRV::OpFOrdEqual;
3112 return SPIRV::OpFOrdGreaterThanEqual;
3114 return SPIRV::OpFOrdGreaterThan;
3116 return SPIRV::OpFOrdLessThanEqual;
3118 return SPIRV::OpFOrdLessThan;
3120 return SPIRV::OpFOrdNotEqual;
3122 return SPIRV::OpOrdered;
3124 return SPIRV::OpFUnordEqual;
3126 return SPIRV::OpFUnordGreaterThanEqual;
3128 return SPIRV::OpFUnordGreaterThan;
3130 return SPIRV::OpFUnordLessThanEqual;
3132 return SPIRV::OpFUnordLessThan;
3134 return SPIRV::OpFUnordNotEqual;
3136 return SPIRV::OpUnordered;
3146 return SPIRV::OpIEqual;
3148 return SPIRV::OpINotEqual;
3150 return SPIRV::OpSGreaterThanEqual;
3152 return SPIRV::OpSGreaterThan;
3154 return SPIRV::OpSLessThanEqual;
3156 return SPIRV::OpSLessThan;
3158 return SPIRV::OpUGreaterThanEqual;
3160 return SPIRV::OpUGreaterThan;
3162 return SPIRV::OpULessThanEqual;
3164 return SPIRV::OpULessThan;
3173 return SPIRV::OpPtrEqual;
3175 return SPIRV::OpPtrNotEqual;
3186 return SPIRV::OpLogicalEqual;
3188 return SPIRV::OpLogicalNotEqual;
3226bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3227 SPIRVTypeInst ResType,
3229 unsigned OpAnyOrAll)
const {
3230 assert(
I.getNumOperands() == 3);
3231 assert(
I.getOperand(2).isReg());
3233 Register InputRegister =
I.getOperand(2).getReg();
3236 assert(InputType &&
"VReg has no type assigned");
3240 assert(ResVReg ==
I.getOperand(0).getReg());
3241 return BuildCOPY(ResVReg, InputRegister,
I);
3245 unsigned SpirvNotEqualId =
3246 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3248 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3253 IsBoolTy ? InputRegister
3261 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3263 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3280bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3281 SPIRVTypeInst ResType,
3282 MachineInstr &
I)
const {
3283 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3286bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3287 SPIRVTypeInst ResType,
3288 MachineInstr &
I)
const {
3289 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3293bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3294 SPIRVTypeInst ResType,
3295 MachineInstr &
I)
const {
3296 assert(
I.getNumOperands() == 4);
3297 assert(
I.getOperand(2).isReg());
3298 assert(
I.getOperand(3).isReg());
3300 [[maybe_unused]] SPIRVTypeInst VecType =
3305 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3306 "dot product requires either a vector of at least 2 components or"
3307 " the SPV_EXT_long vector extension.");
3309 [[maybe_unused]] SPIRVTypeInst EltType =
3318 .
addUse(
I.getOperand(2).getReg())
3319 .
addUse(
I.getOperand(3).getReg())
3324bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3325 SPIRVTypeInst ResType,
3328 assert(
I.getNumOperands() == 4);
3329 assert(
I.getOperand(2).isReg());
3330 assert(
I.getOperand(3).isReg());
3333 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3337 .
addUse(
I.getOperand(2).getReg())
3338 .
addUse(
I.getOperand(3).getReg())
3345bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3346 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3347 assert(
I.getNumOperands() == 4);
3348 assert(
I.getOperand(2).isReg());
3349 assert(
I.getOperand(3).isReg());
3353 Register Vec0 =
I.getOperand(2).getReg();
3354 Register Vec1 =
I.getOperand(3).getReg();
3358 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3367 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3368 "dot product requires either a vector of at least 2 components "
3369 "or the SPV_EXT_long_vector extension.");
3372 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3382 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3393 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3405bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3406 SPIRVTypeInst ResType,
3407 MachineInstr &
I)
const {
3409 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3412 .
addUse(
I.getOperand(2).getReg())
3417bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3418 SPIRVTypeInst ResType,
3419 MachineInstr &
I)
const {
3421 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3424 .
addUse(
I.getOperand(2).getReg())
3429bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3430 SPIRVTypeInst ResType,
3431 MachineInstr &
I)
const {
3434 Register Src =
I.getOperand(2).getReg();
3477bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3478 SPIRVTypeInst ResType,
3479 MachineInstr &
I)
const {
3481 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3484 .
addUse(
I.getOperand(2).getReg())
3489template <
bool Signed>
3490bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3491 SPIRVTypeInst ResType,
3492 MachineInstr &
I)
const {
3493 assert(
I.getNumOperands() == 5);
3494 assert(
I.getOperand(2).isReg());
3495 assert(
I.getOperand(3).isReg());
3496 assert(
I.getOperand(4).isReg());
3499 Register Acc =
I.getOperand(2).getReg();
3503 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3505 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3510 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3513 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3525template <
bool Signed>
3526bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3527 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3528 assert(
I.getNumOperands() == 5);
3529 assert(
I.getOperand(2).isReg());
3530 assert(
I.getOperand(3).isReg());
3531 assert(
I.getOperand(4).isReg());
3534 Register Acc =
I.getOperand(2).getReg();
3540 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3544 for (
unsigned i = 0; i < 4; i++) {
3567 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3587 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3602bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3603 SPIRVTypeInst ResType,
3604 MachineInstr &
I)
const {
3605 assert(
I.getNumOperands() == 3);
3606 assert(
I.getOperand(2).isReg());
3608 Register VZero = buildZerosValF(ResType,
I);
3609 Register VOne = buildOnesValF(ResType,
I);
3611 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3614 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3616 .
addUse(
I.getOperand(2).getReg())
3623bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3624 SPIRVTypeInst ResType,
3625 MachineInstr &
I)
const {
3626 assert(
I.getNumOperands() == 3);
3627 assert(
I.getOperand(2).isReg());
3629 Register InputRegister =
I.getOperand(2).getReg();
3631 auto &
DL =
I.getDebugLoc();
3634 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3641 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3643 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3651 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3656 if (NeedsConversion) {
3657 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3668bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3669 SPIRVTypeInst ResType,
3671 unsigned Opcode)
const {
3675 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3681 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3682 BMI.addUse(
I.getOperand(J).getReg());
3689bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3692 bool WithGroupSync)
const {
3694 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3696 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3698 assert(((Scope != SPIRV::Scope::Workgroup) ||
3699 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3700 "Workgroup Scope must set WorkGroupMemory semantic "
3701 "in Barrier instruction");
3703 assert(((Scope != SPIRV::Scope::Device) ||
3704 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3705 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3706 "Device Scope must set UniformMemory and ImageMemory semantic "
3707 "in Barrier instruction");
3713 if (WithGroupSync) {
3714 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3718 Register ScopeReg = buildI32Constant(Scope,
I);
3719 Register MemSemReg = buildI32Constant(MemSem,
I);
3721 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3725bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3726 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3731 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3732 SPIRV::OpGroupNonUniformBallot))
3737 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3742 .
addImm(SPIRV::GroupOperation::Reduce)
3749bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3750 SPIRVTypeInst ResType,
3751 MachineInstr &
I)
const {
3756 Register InputReg =
I.getOperand(2).getReg();
3761 bool IsVector = NumElems > 1 ||
3762 (InputType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
3776 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3777 SPIRV::OpGroupNonUniformAllEqual);
3782 ElementResults.
reserve(NumElems);
3784 for (
unsigned Idx = 0;
Idx < NumElems; ++
Idx) {
3797 ElemInput = Extracted;
3803 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3814 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3825bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3826 SPIRVTypeInst ResType,
3827 MachineInstr &
I)
const {
3829 assert(
I.getNumOperands() == 3);
3831 auto Op =
I.getOperand(2);
3841 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3843 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3844 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3865 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3869 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3876bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3877 SPIRVTypeInst ResType,
3879 bool IsUnsigned)
const {
3880 return selectWaveReduce(
3881 ResVReg, ResType,
I, IsUnsigned,
3882 [&](
Register InputRegister,
bool IsUnsigned) {
3883 const bool IsFloatTy =
3885 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3886 : SPIRV::OpGroupNonUniformSMax;
3887 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3891bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3892 SPIRVTypeInst ResType,
3894 bool IsUnsigned)
const {
3895 return selectWaveReduce(
3896 ResVReg, ResType,
I, IsUnsigned,
3897 [&](
Register InputRegister,
bool IsUnsigned) {
3898 const bool IsFloatTy =
3900 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3901 : SPIRV::OpGroupNonUniformSMin;
3902 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3906bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3907 SPIRVTypeInst ResType,
3908 MachineInstr &
I)
const {
3909 return selectWaveReduce(ResVReg, ResType,
I,
false,
3910 [&](
Register InputRegister,
bool IsUnsigned) {
3912 InputRegister, SPIRV::OpTypeFloat);
3913 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3914 : SPIRV::OpGroupNonUniformIAdd;
3918bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3919 SPIRVTypeInst ResType,
3920 MachineInstr &
I)
const {
3921 return selectWaveReduce(ResVReg, ResType,
I,
false,
3922 [&](
Register InputRegister,
bool IsUnsigned) {
3924 InputRegister, SPIRV::OpTypeFloat);
3925 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3926 : SPIRV::OpGroupNonUniformIMul;
3930template <
typename PickOpcodeFn>
3931bool SPIRVInstructionSelector::selectWaveReduce(
3932 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3933 PickOpcodeFn &&PickOpcode)
const {
3934 assert(
I.getNumOperands() == 3);
3935 assert(
I.getOperand(2).isReg());
3937 Register InputRegister =
I.getOperand(2).getReg();
3941 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3944 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3950 .
addImm(SPIRV::GroupOperation::Reduce)
3951 .
addUse(
I.getOperand(2).getReg())
3956bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3957 SPIRVTypeInst ResType,
3959 unsigned Opcode)
const {
3960 return selectWaveReduce(
3961 ResVReg, ResType,
I,
false,
3962 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3965bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3966 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3967 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3968 [&](
Register InputRegister,
bool IsUnsigned) {
3970 InputRegister, SPIRV::OpTypeFloat);
3972 ? SPIRV::OpGroupNonUniformFAdd
3973 : SPIRV::OpGroupNonUniformIAdd;
3977bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3978 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3979 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3980 [&](
Register InputRegister,
bool IsUnsigned) {
3982 InputRegister, SPIRV::OpTypeFloat);
3984 ? SPIRV::OpGroupNonUniformFMul
3985 : SPIRV::OpGroupNonUniformIMul;
3989template <
typename PickOpcodeFn>
3990bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3991 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3992 PickOpcodeFn &&PickOpcode)
const {
3993 assert(
I.getNumOperands() == 3);
3994 assert(
I.getOperand(2).isReg());
3996 Register InputRegister =
I.getOperand(2).getReg();
4000 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
4003 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
4009 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
4010 .
addUse(
I.getOperand(2).getReg())
4015bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
4016 SPIRVTypeInst ResType,
4019 assert(
I.getNumOperands() == 3);
4020 assert(
I.getOperand(2).isReg());
4022 Register InputRegister =
I.getOperand(2).getReg();
4028 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
4039bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
4040 SPIRVTypeInst ResType,
4047 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
4052 : SPIRV::OpUConvert;
4054 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4057 ShiftOp = SPIRV::OpShiftRightLogicalV;
4062 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4063 TII.get(SPIRV::OpConstantComposite))
4066 for (
unsigned It = 0; It <
N; ++It)
4070 ShiftConst = CompositeReg;
4075 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
4080 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
4085 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
4090 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
4093bool SPIRVInstructionSelector::handle64BitOverflow(
4095 unsigned int Opcode,
4102 "handle64BitOverflow should only be used for integer types");
4104 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4106 MachineIRBuilder MIRBuilder(
I);
4108 SPIRVTypeInst I64x2Type =
4110 SPIRVTypeInst Vec2ResType =
4113 std::vector<Register> PartialRegs;
4115 unsigned CurrentComponent = 0;
4116 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4120 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4121 TII.get(SPIRV::OpVectorShuffle))
4126 .
addImm(CurrentComponent)
4127 .
addImm(CurrentComponent + 1);
4137 PartialRegs.push_back(SubVecReg);
4140 if (CurrentComponent != ComponentCount) {
4146 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4147 SPIRV::OpVectorExtractDynamic))
4156 PartialRegs.push_back(FinalElemResReg);
4160 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4161 SPIRV::OpCompositeConstruct);
4164bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4165 SPIRVTypeInst ResType,
4169 if (ComponentCount > 2)
4170 return handle64BitOverflow(
4171 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4173 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4175 MachineIRBuilder MIRBuilder(
I);
4179 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4183 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4188 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4195 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4196 TII.get(SPIRV::OpVectorShuffle))
4201 for (
unsigned J = 0; J < ComponentCount; ++J) {
4208 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4211bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4212 SPIRVTypeInst ResType,
4216 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4224bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4225 SPIRVTypeInst ResType,
4226 MachineInstr &
I)
const {
4227 Register OpReg =
I.getOperand(1).getReg();
4236 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4238 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4240 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4242 return SPIRVInstructionSelector::diagnoseUnsupported(
4243 I,
"G_BITREVERSE only support 16,32,64 bits.");
4247 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4258 unsigned AndOp = SPIRV::OpBitwiseAndS;
4259 unsigned OrOp = SPIRV::OpBitwiseOrS;
4260 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4261 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4262 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4264 AndOp = SPIRV::OpBitwiseAndV;
4265 OrOp = SPIRV::OpBitwiseOrV;
4266 ShlOp = SPIRV::OpShiftLeftLogicalV;
4267 ShrOp = SPIRV::OpShiftRightLogicalV;
4273 const unsigned Shift) ->
Register {
4276 (ResType->
getOpcode() != SPIRV::OpTypeVectorIdEXT ||
4283 Register MaskReg = CreateConst(Mask);
4284 Register ShiftReg = CreateConst(Shift);
4291 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4292 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4293 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4294 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4295 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4304 while ((Shift >>= 1) > 0) {
4311 return BuildCOPY(ResVReg, Result,
I);
4314bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4315 SPIRVTypeInst ResType,
4316 MachineInstr &
I)
const {
4317 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4318 "G_FREEZE must define and use a register");
4319 Register OpReg =
I.getOperand(1).getReg();
4323 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4336 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4337 if (
Def->getOpcode() == TargetOpcode::COPY)
4340 switch (
Def->getOpcode()) {
4341 case SPIRV::ASSIGN_TYPE:
4342 if (MachineInstr *AssignToDef =
4344 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4345 Reg =
Def->getOperand(2).getReg();
4348 case SPIRV::OpUndef:
4349 Reg =
Def->getOperand(1).getReg();
4352 unsigned DestOpCode;
4354 DestOpCode = SPIRV::OpConstantNull;
4355 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4356 "static undef/poison lowered to OpConstantNull\n");
4358 DestOpCode = TargetOpcode::COPY;
4360 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4361 "skipped, lowered as a copy of the operand\n");
4363 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4364 .
addDef(
I.getOperand(0).getReg())
4372bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4373 SPIRVTypeInst ResType,
4374 MachineInstr &
I)
const {
4378 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4382 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4387 for (
unsigned i =
I.getNumExplicitDefs();
4388 i <
I.getNumExplicitOperands() && IsConst; ++i)
4393 return diagnoseUnsupported(
4394 I,
"There must be at least two constituent operands in a vector");
4399 for (
unsigned i =
I.getNumExplicitDefs();
4400 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4401 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4406 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4413 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4414 TII.get(IsConst ? SPIRV::OpConstantComposite
4415 : SPIRV::OpCompositeConstruct))
4418 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4419 MIB.
addUse(
I.getOperand(i).getReg());
4424bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4425 SPIRVTypeInst ResType,
4426 MachineInstr &
I)
const {
4430 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4435 unsigned OpIdx =
I.getNumExplicitDefs();
4436 if (!
I.getOperand(OpIdx).isReg())
4440 Register OpReg =
I.getOperand(OpIdx).getReg();
4444 return diagnoseUnsupported(
4445 I,
"There must be at least two constituent operands in a vector");
4448 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4449 TII.get(IsConst ? SPIRV::OpConstantComposite
4450 : SPIRV::OpCompositeConstruct))
4453 for (
unsigned i = 0; i <
N; ++i)
4459bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4460 SPIRVTypeInst ResType,
4461 MachineInstr &
I)
const {
4467 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4469 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4470 TII.get(SPIRV::OpCompositeConstruct))
4473 for (
unsigned OpIdx =
I.getNumExplicitDefs();
4474 OpIdx <
I.getNumExplicitOperands(); ++OpIdx)
4475 MIB.
addUse(
I.getOperand(OpIdx).getReg());
4480bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4481 SPIRVTypeInst ResType,
4482 MachineInstr &
I)
const {
4488 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4490 Opcode = SPIRV::OpDemoteToHelperInvocation;
4492 Opcode = SPIRV::OpKill;
4497 ToErase.eraseFromParent();
4506bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4507 SPIRVTypeInst ResType,
unsigned CmpOpc,
4508 MachineInstr &
I)
const {
4509 Register Cmp0 =
I.getOperand(2).getReg();
4510 Register Cmp1 =
I.getOperand(3).getReg();
4513 "CMP operands should have the same type");
4514 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4524bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4525 SPIRVTypeInst ResType,
4526 MachineInstr &
I)
const {
4527 auto Pred =
I.getOperand(1).getPredicate();
4530 Register CmpOperand =
I.getOperand(2).getReg();
4532 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4537 Register Op1 =
I.getOperand(3).getReg();
4541 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4546 I.getOperand(3).setReg(NewOp1);
4552 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4556SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4557 SPIRVTypeInst ResType)
const {
4559 SPIRVTypeInst SpvI32Ty =
4562 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4569 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4572 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4575 .
addImm(APInt(32, Val).getZExtValue());
4577 GR.
add(ConstInt,
MI);
4584Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4585 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4587 SPIRVTypeInst SpvI32Ty =
4589 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4594 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4595 MachineInstr *
MI =
nullptr;
4599 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4603 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4604 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4610 GR.
add(ConstInt,
MI);
4615bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4616 SPIRVTypeInst ResType,
4617 MachineInstr &
I)
const {
4619 return selectCmp(ResVReg, ResType, CmpOp,
I);
4622bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4623 SPIRVTypeInst ResType,
4624 MachineInstr &
I)
const {
4626 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4636 MachineIRBuilder MIRBuilder(
I);
4643 APFloat ConstVal(3.3219280948873623);
4647 APFloat::rmNearestTiesToEven, &LosesInfo);
4652 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
4654 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4655 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4657 if (!selectExtInst(ResVReg, ResType,
I,
4658 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4668Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4669 MachineInstr &
I)
const {
4677bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4683 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4691 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4694 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4695 Def->getOpcode() == SPIRV::OpConstantI)
4708 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4709 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4711 Intrinsic::spv_const_composite)) {
4712 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4713 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4714 if (!IsZero(
Def->getOperand(i).getReg()))
4723Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4724 MachineInstr &
I)
const {
4733Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4734 MachineInstr &
I)
const {
4744 SPIRVTypeInst ResType,
4745 MachineInstr &
I)
const {
4756 MachineInstr &
I)
const {
4761 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4762 TII.get(SPIRV::OpCompositeConstruct))
4765 for (
unsigned J = 0; J < NumElts; ++J)
4771bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4772 SPIRVTypeInst ResType,
4773 MachineInstr &
I)
const {
4774 Register SelectFirstArg =
I.getOperand(2).getReg();
4775 Register SelectSecondArg =
I.getOperand(3).getReg();
4784 Register CondReg =
I.getOperand(1).getReg();
4785 bool IsScalarBool = GR.
isScalarOfType(CondReg, SPIRV::OpTypeBool);
4793 CondReg = buildVectorSplat(CondReg, NumElts,
I);
4794 IsScalarBool =
false;
4797 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4798 }
else if (IsPtrTy) {
4799 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4801 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4804 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4805 "boolean condition");
4807 Opcode = SPIRV::OpSelectSFSCond;
4808 }
else if (IsPtrTy) {
4809 Opcode = SPIRV::OpSelectSPSCond;
4811 Opcode = SPIRV::OpSelectSISCond;
4814 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4826bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4827 SPIRVTypeInst ResType,
4829 MachineInstr &InsertAt,
4830 bool IsSigned)
const {
4832 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4833 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4834 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4836 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4848bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4849 SPIRVTypeInst ResType,
4850 MachineInstr &
I,
bool IsSigned,
4851 unsigned Opcode)
const {
4852 Register SrcReg =
I.getOperand(1).getReg();
4863 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4865 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4868bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4869 SPIRVTypeInst ResType, MachineInstr &
I,
4870 bool IsSigned)
const {
4871 Register SrcReg =
I.getOperand(1).getReg();
4873 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4877 if (ResType == SrcType)
4878 return BuildCOPY(ResVReg, SrcReg,
I);
4880 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4881 return selectUnOp(ResVReg, ResType,
I, Opcode);
4884bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4885 SPIRVTypeInst ResType,
4887 bool IsSigned)
const {
4888 MachineIRBuilder MIRBuilder(
I);
4889 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
4894 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4902 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4905 .
addUse(
I.getOperand(1).getReg())
4906 .
addUse(
I.getOperand(2).getReg())
4911 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4914 .
addUse(
I.getOperand(1).getReg())
4915 .
addUse(
I.getOperand(2).getReg())
4923 unsigned SelectOpcode =
4924 (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4926 ? SPIRV::OpSelectVIVCond
4927 : SPIRV::OpSelectSISCond;
4932 .
addUse(buildOnesVal(
true, ResType,
I))
4933 .
addUse(buildZerosVal(ResType,
I))
4940 .
addUse(buildOnesVal(
false, ResType,
I))
4945bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4948 SPIRVTypeInst IntTy,
4949 SPIRVTypeInst BoolTy)
const {
4953 isVectorType(IntTy) ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4955 Register One = buildOnesVal(
false, IntTy,
I);
4963 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4972bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4973 SPIRVTypeInst ResType,
4974 MachineInstr &
I)
const {
4975 Register IntReg =
I.getOperand(1).getReg();
4978 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4979 if (ArgType == ResType)
4980 return BuildCOPY(ResVReg, IntReg,
I);
4982 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4983 return selectUnOp(ResVReg, ResType,
I, Opcode);
4986bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4987 SPIRVTypeInst ResType,
4988 MachineInstr &
I)
const {
4989 unsigned Opcode =
I.getOpcode();
4990 unsigned TpOpcode = ResType->
getOpcode();
4992 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
4993 assert(Opcode == TargetOpcode::G_CONSTANT &&
4994 I.getOperand(1).getCImm()->isZero());
4995 MachineBasicBlock &DepMBB =
I.getMF()->front();
4998 }
else if (TpOpcode == SPIRV::OpTypeVectorIdEXT) {
5003 "Expected <1 x T> Vector!");
5004 if (Opcode == TargetOpcode::G_FCONSTANT)
5010 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
5018 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
5021bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
5022 SPIRVTypeInst ResType,
5023 MachineInstr &
I)
const {
5024 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5031bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
5032 SPIRVTypeInst ResType,
5033 MachineInstr &
I)
const {
5035 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
5039 .
addUse(
I.getOperand(3).getReg())
5041 .
addUse(
I.getOperand(2).getReg());
5042 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
5048bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
5049 SPIRVTypeInst ResType,
5050 MachineInstr &
I)
const {
5051 Type *MaybeResTy =
nullptr;
5056 "Expected aggregate type for extractv instruction");
5058 SPIRV::AccessQualifier::ReadWrite,
false);
5062 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5065 .
addUse(
I.getOperand(2).getReg());
5066 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
5072bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
5073 SPIRVTypeInst ResType,
5074 MachineInstr &
I)
const {
5075 if (
getImm(
I.getOperand(4), MRI))
5076 return selectInsertVal(ResVReg, ResType,
I);
5078 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
5081 .
addUse(
I.getOperand(2).getReg())
5082 .
addUse(
I.getOperand(3).getReg())
5083 .
addUse(
I.getOperand(4).getReg())
5088bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
5089 SPIRVTypeInst ResType,
5090 MachineInstr &
I)
const {
5091 if (
getImm(
I.getOperand(3), MRI))
5092 return selectExtractVal(ResVReg, ResType,
I);
5094 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
5097 .
addUse(
I.getOperand(2).getReg())
5098 .
addUse(
I.getOperand(3).getReg())
5103bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
5104 SPIRVTypeInst ResType,
5105 MachineInstr &
I)
const {
5106 const bool IsGEPInBounds =
I.getOperand(2).getImm();
5109 const bool UseUntypedPointers =
5110 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
5115 if (UseUntypedPointers) {
5117 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
5118 : SPIRV::OpUntypedAccessChainKHR;
5120 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
5121 : SPIRV::OpUntypedPtrAccessChainKHR;
5130 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
5132 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
5133 : SPIRV::OpPtrAccessChain;
5138 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5143 if (UseUntypedPointers) {
5158 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5159 Def->getOperand(1).isReg())
5161 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5162 if (
const auto *GVar =
5165 SPIRV::AccessQualifier::ReadWrite,
5169 return diagnoseUnsupported(
5170 I,
"could not deduce the base type of an untyped access chain");
5175 Res.addUse(BaseReg);
5177 const bool IsAccessChainOpcode =
5178 (Opcode == SPIRV::OpAccessChain ||
5179 Opcode == SPIRV::OpInBoundsAccessChain ||
5180 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5181 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5183 assert((!IsAccessChainOpcode || (
getImm(
I.getOperand(4), MRI) &&
5184 foldImm(
I.getOperand(4), MRI) == 0)) &&
5185 "Cannot translate GEP to OpAccessChain.");
5188 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5189 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5190 Res.addUse(
I.getOperand(i).getReg());
5191 Res.constrainAllUses(
TII,
TRI, RBI);
5200 if (Extract.
getOpcode() == SPIRV::OpCompositeExtract) {
5206 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5212 TII.get(SPIRV::OpCompositeInsert))
5226bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5228 unsigned Lim =
I.getNumExplicitOperands();
5229 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5230 Register OpReg =
I.getOperand(i).getReg();
5231 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
5233 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
5234 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5235 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5248 SPIRVTypeInst WrapType = OpType;
5249 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5251 SPIRV::StorageClass::CodeSectionINTEL) {
5253 SPIRV::StorageClass::Function,
I);
5260 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5261 TII.get(SPIRV::OpSpecConstantOp))
5264 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5266 GR.
add(OpDefine, MIB);
5272bool SPIRVInstructionSelector::selectDerivativeInst(
5273 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5274 const unsigned DPdOpCode)
const {
5277 if (!errorIfInstrOutsideShader(
I))
5283 Register SrcReg =
I.getOperand(2).getReg();
5288 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5291 .
addUse(
I.getOperand(2).getReg());
5293 MachineIRBuilder MIRBuilder(
I);
5296 if (componentCount != 1)
5304 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5309 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5314 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5322bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5323 SPIRVTypeInst ResType,
5324 MachineInstr &
I)
const {
5328 case Intrinsic::spv_load:
5329 return selectLoad(ResVReg, ResType,
I);
5330 case Intrinsic::spv_atomic_load:
5331 return selectAtomicLoad(ResVReg, ResType,
I);
5332 case Intrinsic::spv_store:
5333 return selectStore(
I);
5334 case Intrinsic::spv_atomic_store:
5335 return selectAtomicStore(
I);
5336 case Intrinsic::spv_extractv:
5337 return selectExtractVal(ResVReg, ResType,
I);
5338 case Intrinsic::spv_insertv:
5339 return selectInsertVal(ResVReg, ResType,
I);
5340 case Intrinsic::spv_extractelt:
5341 return selectExtractElt(ResVReg, ResType,
I);
5342 case Intrinsic::spv_insertelt:
5343 return selectInsertElt(ResVReg, ResType,
I);
5344 case Intrinsic::spv_gep:
5345 return selectGEP(ResVReg, ResType,
I);
5346 case Intrinsic::spv_bitcast: {
5347 Register OpReg =
I.getOperand(2).getReg();
5348 SPIRVTypeInst OpType =
5352 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5354 case Intrinsic::spv_unref_global:
5355 case Intrinsic::spv_init_global: {
5356 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
5361 Register GVarVReg =
MI->getOperand(0).getReg();
5362 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5367 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5369 MI->eraseFromParent();
5373 case Intrinsic::spv_undef: {
5374 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5380 case Intrinsic::spv_poison:
5381 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5386 case Intrinsic::spv_freeze:
5387 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5390 .
addUse(
I.getOperand(2).getReg())
5393 case Intrinsic::spv_named_boolean_spec_constant: {
5394 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5395 : SPIRV::OpSpecConstantFalse;
5397 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5398 .
addDef(
I.getOperand(0).getReg())
5401 unsigned SpecId =
I.getOperand(2).getImm();
5403 SPIRV::Decoration::SpecId, {SpecId});
5407 case Intrinsic::spv_const_composite: {
5409 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5415 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5417 std::function<bool(
Register)> HasSpecConstOperand =
5427 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5428 J < Def->getNumExplicitOperands(); ++J) {
5429 if (
Def->getOperand(J).isReg() &&
5430 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5436 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5437 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5438 : SPIRV::OpConstantComposite;
5439 unsigned ContinuedOpc = HasSpecConst
5440 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5441 : SPIRV::OpConstantCompositeContinuedINTEL;
5442 MachineIRBuilder MIR(
I);
5444 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5446 for (
auto *Instr : Instructions) {
5447 Instr->setDebugLoc(
I.getDebugLoc());
5452 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5459 case Intrinsic::spv_assign_name: {
5460 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5461 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5462 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5463 i <
I.getNumExplicitOperands(); ++i) {
5464 MIB.
addImm(
I.getOperand(i).getImm());
5469 case Intrinsic::spv_switch: {
5470 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5471 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5472 if (
I.getOperand(i).isReg())
5473 MIB.
addReg(
I.getOperand(i).getReg());
5474 else if (
I.getOperand(i).isCImm())
5475 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5476 else if (
I.getOperand(i).isMBB())
5477 MIB.
addMBB(
I.getOperand(i).getMBB());
5484 case Intrinsic::spv_loop_merge: {
5485 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5486 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5487 if (
I.getOperand(i).isMBB())
5488 MIB.
addMBB(
I.getOperand(i).getMBB());
5495 case Intrinsic::spv_loop_control_intel: {
5497 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5498 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5503 case Intrinsic::spv_selection_merge: {
5505 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5506 assert(
I.getOperand(1).isMBB() &&
5507 "operand 1 to spv_selection_merge must be a basic block");
5508 MIB.
addMBB(
I.getOperand(1).getMBB());
5509 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5513 case Intrinsic::spv_cmpxchg:
5514 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5515 case Intrinsic::spv_unreachable:
5516 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5519 case Intrinsic::spv_abort:
5520 return selectAbort(
I);
5521 case Intrinsic::spv_alloca:
5522 return selectFrameIndex(ResVReg, ResType,
I);
5523 case Intrinsic::spv_alloca_array:
5524 return selectAllocaArray(ResVReg, ResType,
I);
5525 case Intrinsic::spv_assume:
5527 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5528 .
addUse(
I.getOperand(1).getReg())
5533 case Intrinsic::spv_expect:
5535 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5538 .
addUse(
I.getOperand(2).getReg())
5539 .
addUse(
I.getOperand(3).getReg())
5544 case Intrinsic::arithmetic_fence:
5545 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5546 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5549 .
addUse(
I.getOperand(2).getReg())
5553 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5555 case Intrinsic::spv_thread_id:
5561 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5563 case Intrinsic::spv_thread_id_in_group:
5569 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5571 case Intrinsic::spv_group_id:
5577 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5579 case Intrinsic::spv_flattened_thread_id_in_group:
5586 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5588 case Intrinsic::spv_workgroup_size:
5589 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5591 case Intrinsic::spv_global_size:
5592 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5594 case Intrinsic::spv_global_offset:
5595 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5597 case Intrinsic::spv_num_workgroups:
5598 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5600 case Intrinsic::spv_subgroup_size:
5601 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5603 case Intrinsic::spv_num_subgroups:
5604 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5606 case Intrinsic::spv_subgroup_id:
5607 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5608 case Intrinsic::spv_subgroup_local_invocation_id:
5609 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5610 ResVReg, ResType,
I);
5611 case Intrinsic::spv_subgroup_max_size:
5612 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5614 case Intrinsic::spv_fdot:
5615 return selectFloatDot(ResVReg, ResType,
I);
5616 case Intrinsic::spv_udot:
5617 case Intrinsic::spv_sdot:
5618 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5620 return selectIntegerDot(ResVReg, ResType,
I,
5621 IID == Intrinsic::spv_sdot);
5622 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5623 case Intrinsic::spv_dot4add_i8packed:
5624 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5626 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5627 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5628 case Intrinsic::spv_dot4add_u8packed:
5629 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5631 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5632 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5633 case Intrinsic::spv_all:
5634 return selectAll(ResVReg, ResType,
I);
5635 case Intrinsic::spv_any:
5636 return selectAny(ResVReg, ResType,
I);
5637 case Intrinsic::spv_distance:
5638 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5639 case Intrinsic::spv_lerp:
5640 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5641 case Intrinsic::spv_length:
5642 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5643 case Intrinsic::spv_degrees:
5644 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5645 case Intrinsic::spv_faceforward:
5646 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5647 case Intrinsic::spv_frac:
5648 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5649 case Intrinsic::spv_isinf:
5650 return selectOpIsInf(ResVReg, ResType,
I);
5651 case Intrinsic::spv_isnan:
5652 return selectOpIsNan(ResVReg, ResType,
I);
5653 case Intrinsic::spv_isfinite:
5654 return selectOpIsFinite(ResVReg, ResType,
I);
5655 case Intrinsic::spv_isnormal:
5656 return selectOpIsNormal(ResVReg, ResType,
I);
5657 case Intrinsic::spv_normalize:
5658 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5659 case Intrinsic::spv_refract:
5660 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5661 case Intrinsic::spv_reflect:
5662 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5663 case Intrinsic::spv_rsqrt:
5664 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5665 case Intrinsic::spv_sign:
5666 return selectSign(ResVReg, ResType,
I);
5667 case Intrinsic::spv_smoothstep:
5668 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5669 case Intrinsic::spv_firstbituhigh:
5670 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5671 case Intrinsic::spv_firstbitshigh:
5672 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5673 case Intrinsic::spv_firstbitlow:
5674 return selectFirstBitLow(ResVReg, ResType,
I);
5675 case Intrinsic::spv_all_memory_barrier:
5676 return selectBarrierInst(
I, SPIRV::Scope::Device,
5677 SPIRV::MemorySemantics::UniformMemory |
5678 SPIRV::MemorySemantics::ImageMemory |
5679 SPIRV::MemorySemantics::WorkgroupMemory,
5681 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5682 return selectBarrierInst(
I, SPIRV::Scope::Device,
5683 SPIRV::MemorySemantics::UniformMemory |
5684 SPIRV::MemorySemantics::ImageMemory |
5685 SPIRV::MemorySemantics::WorkgroupMemory,
5687 case Intrinsic::spv_device_memory_barrier:
5688 return selectBarrierInst(
I, SPIRV::Scope::Device,
5689 SPIRV::MemorySemantics::UniformMemory |
5690 SPIRV::MemorySemantics::ImageMemory,
5692 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5693 return selectBarrierInst(
I, SPIRV::Scope::Device,
5694 SPIRV::MemorySemantics::UniformMemory |
5695 SPIRV::MemorySemantics::ImageMemory,
5697 case Intrinsic::spv_group_memory_barrier:
5698 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5699 SPIRV::MemorySemantics::WorkgroupMemory,
5701 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5702 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5703 SPIRV::MemorySemantics::WorkgroupMemory,
5705 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5706 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5707 SPIRV::StorageClass::StorageClass ResSC =
5710 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5711 "from the Generic storage class");
5712 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5720 case Intrinsic::spv_lifetime_start:
5721 case Intrinsic::spv_lifetime_end: {
5722 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5723 : SPIRV::OpLifetimeStop;
5724 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5725 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5734 case Intrinsic::spv_saturate:
5735 return selectSaturate(ResVReg, ResType,
I);
5736 case Intrinsic::spv_nclamp:
5737 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5738 case Intrinsic::spv_uclamp:
5739 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5740 case Intrinsic::spv_sclamp:
5741 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5742 case Intrinsic::spv_subgroup_prefix_bit_count:
5743 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5744 case Intrinsic::spv_wave_active_countbits:
5745 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5746 case Intrinsic::spv_wave_all_equal:
5747 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5748 case Intrinsic::spv_wave_all:
5749 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5750 case Intrinsic::spv_wave_any:
5751 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5752 case Intrinsic::spv_subgroup_ballot:
5753 return selectWaveOpInst(ResVReg, ResType,
I,
5754 SPIRV::OpGroupNonUniformBallot);
5755 case Intrinsic::spv_wave_is_first_lane:
5756 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5757 case Intrinsic::spv_wave_reduce_or:
5758 return selectWaveReduceOp(ResVReg, ResType,
I,
5759 SPIRV::OpGroupNonUniformBitwiseOr);
5760 case Intrinsic::spv_wave_reduce_xor:
5761 return selectWaveReduceOp(ResVReg, ResType,
I,
5762 SPIRV::OpGroupNonUniformBitwiseXor);
5763 case Intrinsic::spv_wave_reduce_and:
5764 return selectWaveReduceOp(ResVReg, ResType,
I,
5765 SPIRV::OpGroupNonUniformBitwiseAnd);
5766 case Intrinsic::spv_wave_reduce_umax:
5767 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5768 case Intrinsic::spv_wave_reduce_max:
5769 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5770 case Intrinsic::spv_wave_reduce_umin:
5771 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5772 case Intrinsic::spv_wave_reduce_min:
5773 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5774 case Intrinsic::spv_wave_reduce_sum:
5775 return selectWaveReduceSum(ResVReg, ResType,
I);
5776 case Intrinsic::spv_wave_product:
5777 return selectWaveReduceProduct(ResVReg, ResType,
I);
5778 case Intrinsic::spv_wave_readlane:
5779 return selectWaveOpInst(ResVReg, ResType,
I,
5780 SPIRV::OpGroupNonUniformShuffle);
5781 case Intrinsic::spv_wave_readlane_first:
5782 return selectWaveOpInst(ResVReg, ResType,
I,
5783 SPIRV::OpGroupNonUniformBroadcastFirst);
5784 case Intrinsic::spv_wave_prefix_sum:
5785 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5786 case Intrinsic::spv_wave_prefix_product:
5787 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5788 case Intrinsic::spv_quad_read_across_x: {
5789 return selectQuadSwap(ResVReg, ResType,
I, 0);
5791 case Intrinsic::spv_quad_read_across_y: {
5792 return selectQuadSwap(ResVReg, ResType,
I, 1);
5794 case Intrinsic::spv_quad_read_across_diagonal: {
5795 return selectQuadSwap(ResVReg, ResType,
I, 2);
5797 case Intrinsic::spv_radians:
5798 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5802 case Intrinsic::instrprof_increment:
5803 case Intrinsic::instrprof_increment_step:
5804 case Intrinsic::instrprof_value_profile:
5807 case Intrinsic::spv_value_md:
5809 case Intrinsic::spv_resource_handlefrombinding: {
5810 return selectHandleFromBinding(ResVReg, ResType,
I);
5812 case Intrinsic::spv_resource_counterhandlefrombinding:
5813 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5814 case Intrinsic::spv_resource_updatecounter:
5815 return selectUpdateCounter(ResVReg, ResType,
I);
5816 case Intrinsic::spv_resource_store_typedbuffer: {
5817 return selectImageWriteIntrinsic(
I);
5819 case Intrinsic::spv_resource_load_typedbuffer: {
5820 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5822 case Intrinsic::spv_resource_load_level: {
5823 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5825 case Intrinsic::spv_resource_getdimensions_x:
5826 case Intrinsic::spv_resource_getdimensions_xy:
5827 case Intrinsic::spv_resource_getdimensions_xyz: {
5828 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5830 case Intrinsic::spv_resource_getdimensions_levels_x:
5831 case Intrinsic::spv_resource_getdimensions_levels_xy:
5832 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5833 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5835 case Intrinsic::spv_resource_getdimensions_ms_xy:
5836 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5837 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5839 case Intrinsic::spv_resource_calculate_lod:
5840 case Intrinsic::spv_resource_calculate_lod_unclamped:
5841 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5842 case Intrinsic::spv_resource_sample:
5843 case Intrinsic::spv_resource_sample_clamp:
5844 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5845 case Intrinsic::spv_resource_samplebias:
5846 case Intrinsic::spv_resource_samplebias_clamp:
5847 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5848 case Intrinsic::spv_resource_samplegrad:
5849 case Intrinsic::spv_resource_samplegrad_clamp:
5850 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5851 case Intrinsic::spv_resource_samplelevel:
5852 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5853 case Intrinsic::spv_resource_samplecmp:
5854 case Intrinsic::spv_resource_samplecmp_clamp:
5855 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5856 case Intrinsic::spv_resource_samplecmplevelzero:
5857 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5858 case Intrinsic::spv_resource_gather:
5859 case Intrinsic::spv_resource_gather_cmp:
5860 return selectGatherIntrinsic(ResVReg, ResType,
I);
5861 case Intrinsic::spv_resource_getbasepointer:
5862 case Intrinsic::spv_resource_getpointer: {
5863 return selectResourceGetPointer(ResVReg, ResType,
I);
5865 case Intrinsic::spv_pushconstant_getpointer: {
5866 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5868 case Intrinsic::spv_discard: {
5869 return selectDiscard(ResVReg, ResType,
I);
5871 case Intrinsic::spv_resource_nonuniformindex: {
5872 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5874 case Intrinsic::spv_unpackhalf2x16: {
5875 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5877 case Intrinsic::spv_packhalf2x16: {
5878 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5880 case Intrinsic::spv_ddx:
5881 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5882 case Intrinsic::spv_ddy:
5883 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5884 case Intrinsic::spv_ddx_coarse:
5885 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5886 case Intrinsic::spv_ddy_coarse:
5887 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5888 case Intrinsic::spv_ddx_fine:
5889 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5890 case Intrinsic::spv_ddy_fine:
5891 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5892 case Intrinsic::spv_fwidth:
5893 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5894 case Intrinsic::spv_masked_gather:
5895 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5896 return selectMaskedGather(ResVReg, ResType,
I);
5897 return diagnoseUnsupported(
5898 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5899 case Intrinsic::spv_masked_scatter:
5900 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5901 return selectMaskedScatter(
I);
5902 return diagnoseUnsupported(
5903 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5904 case Intrinsic::returnaddress:
5905 case Intrinsic::frameaddress: {
5907 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5914 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5919bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5920 SPIRVTypeInst ResType,
5921 MachineInstr &
I)
const {
5924 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5931bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5932 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5934 assert(Intr.getIntrinsicID() ==
5935 Intrinsic::spv_resource_counterhandlefrombinding);
5938 Register MainHandleReg = Intr.getOperand(2).getReg();
5940 assert(MainHandleDef->getIntrinsicID() ==
5941 Intrinsic::spv_resource_handlefrombinding);
5945 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5946 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5947 std::string CounterName =
5952 MachineIRBuilder MIRBuilder(
I);
5954 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5956 ArraySize, IndexReg, CounterName, MIRBuilder);
5958 return BuildCOPY(ResVReg, CounterVarReg,
I);
5961bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5962 SPIRVTypeInst ResType,
5963 MachineInstr &
I)
const {
5965 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5967 Register CounterHandleReg = Intr.getOperand(2).getReg();
5968 Register IncrReg = Intr.getOperand(3).getReg();
5975 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5976 assert(CounterVarPointeeType &&
5977 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5978 "Counter variable must be a struct");
5980 SPIRV::StorageClass::StorageBuffer &&
5981 "Counter variable must be in the storage buffer storage class");
5983 "Counter variable must have exactly 1 member in the struct");
5984 const SPIRVTypeInst MemberType =
5987 "Counter variable struct must have a single i32 member");
5991 MachineIRBuilder MIRBuilder(
I);
5993 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5996 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
6002 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6005 .
addUse(CounterHandleReg)
6012 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
6015 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
6018 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
6027 return BuildCOPY(ResVReg, AtomicRes,
I);
6035 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
6043bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
6044 SPIRVTypeInst ResType,
6045 MachineInstr &
I)
const {
6053 Register ImageReg =
I.getOperand(2).getReg();
6061 Register IdxReg =
I.getOperand(3).getReg();
6063 MachineInstr &Pos =
I;
6065 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
6069bool SPIRVInstructionSelector::generateSampleImage(
6072 DebugLoc Loc, MachineInstr &Pos)
const {
6083 if (!loadHandleBeforePosition(NewSamplerReg,
6089 MachineIRBuilder MIRBuilder(Pos);
6102 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
6103 ImOps.Lod.has_value();
6104 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
6105 : SPIRV::OpImageSampleImplicitLod;
6107 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
6108 : SPIRV::OpImageSampleDrefImplicitLod;
6117 MIB.
addUse(*ImOps.Compare);
6119 uint32_t ImageOperands = 0;
6121 ImageOperands |= SPIRV::ImageOperand::Bias;
6123 ImageOperands |= SPIRV::ImageOperand::Lod;
6124 if (ImOps.GradX && ImOps.GradY)
6125 ImageOperands |= SPIRV::ImageOperand::Grad;
6126 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
6128 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6131 "Non-constant offsets are not supported in sample instructions.");
6136 ImageOperands |= SPIRV::ImageOperand::MinLod;
6138 if (ImageOperands != 0) {
6139 MIB.
addImm(ImageOperands);
6140 if (ImageOperands & SPIRV::ImageOperand::Bias)
6142 if (ImageOperands & SPIRV::ImageOperand::Lod)
6144 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6145 MIB.
addUse(*ImOps.GradX);
6146 MIB.
addUse(*ImOps.GradY);
6149 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6150 MIB.
addUse(*ImOps.Offset);
6151 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6152 MIB.
addUse(*ImOps.MinLod);
6159bool SPIRVInstructionSelector::selectImageQuerySize(
6161 std::optional<Register> LodReg)
const {
6163 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6166 "ImageReg is not an image type.");
6168 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6170 unsigned NumComponents = 0;
6172 case SPIRV::Dim::DIM_1D:
6173 case SPIRV::Dim::DIM_Buffer:
6174 NumComponents =
IsArray ? 2 : 1;
6176 case SPIRV::Dim::DIM_2D:
6177 case SPIRV::Dim::DIM_Cube:
6178 case SPIRV::Dim::DIM_Rect:
6179 NumComponents =
IsArray ? 3 : 2;
6181 case SPIRV::Dim::DIM_3D:
6185 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6190 SPIRVTypeInst ResType =
6195 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6205bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6206 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6207 Register ImageReg =
I.getOperand(2).getReg();
6214 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6217bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6218 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6219 Register ImageReg =
I.getOperand(2).getReg();
6228 Register LodReg =
I.getOperand(3).getReg();
6231 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6233 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6240 TII.get(SPIRV::OpImageQueryLevels))
6247 TII.get(SPIRV::OpCompositeConstruct))
6257bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6258 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6259 Register ImageReg =
I.getOperand(2).getReg();
6270 "OpImageQuerySamples requires a multisampled image");
6272 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6280 TII.get(SPIRV::OpImageQuerySamples))
6287 TII.get(SPIRV::OpCompositeConstruct))
6297bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6298 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6299 Register ImageReg =
I.getOperand(2).getReg();
6300 Register SamplerReg =
I.getOperand(3).getReg();
6301 Register CoordinateReg =
I.getOperand(4).getReg();
6317 if (!loadHandleBeforePosition(
6322 MachineIRBuilder MIRBuilder(
I);
6328 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6338 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6345 unsigned ExtractedIndex =
6347 Intrinsic::spv_resource_calculate_lod_unclamped
6351 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6352 TII.get(SPIRV::OpCompositeExtract))
6362bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6363 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6364 Register ImageReg =
I.getOperand(2).getReg();
6365 Register SamplerReg =
I.getOperand(3).getReg();
6366 Register CoordinateReg =
I.getOperand(4).getReg();
6367 ImageOperands ImOps;
6368 if (
I.getNumOperands() > 5)
6369 ImOps.Offset =
I.getOperand(5).getReg();
6370 if (
I.getNumOperands() > 6)
6371 ImOps.MinLod =
I.getOperand(6).getReg();
6372 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6373 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6376bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6377 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6378 Register ImageReg =
I.getOperand(2).getReg();
6379 Register SamplerReg =
I.getOperand(3).getReg();
6380 Register CoordinateReg =
I.getOperand(4).getReg();
6381 ImageOperands ImOps;
6382 ImOps.Bias =
I.getOperand(5).getReg();
6383 if (
I.getNumOperands() > 6)
6384 ImOps.Offset =
I.getOperand(6).getReg();
6385 if (
I.getNumOperands() > 7)
6386 ImOps.MinLod =
I.getOperand(7).getReg();
6387 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6388 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6391bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6392 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6393 Register ImageReg =
I.getOperand(2).getReg();
6394 Register SamplerReg =
I.getOperand(3).getReg();
6395 Register CoordinateReg =
I.getOperand(4).getReg();
6396 ImageOperands ImOps;
6397 ImOps.GradX =
I.getOperand(5).getReg();
6398 ImOps.GradY =
I.getOperand(6).getReg();
6399 if (
I.getNumOperands() > 7)
6400 ImOps.Offset =
I.getOperand(7).getReg();
6401 if (
I.getNumOperands() > 8)
6402 ImOps.MinLod =
I.getOperand(8).getReg();
6403 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6404 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6407bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6408 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6409 Register ImageReg =
I.getOperand(2).getReg();
6410 Register SamplerReg =
I.getOperand(3).getReg();
6411 Register CoordinateReg =
I.getOperand(4).getReg();
6412 ImageOperands ImOps;
6413 ImOps.Lod =
I.getOperand(5).getReg();
6414 if (
I.getNumOperands() > 6)
6415 ImOps.Offset =
I.getOperand(6).getReg();
6416 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6417 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6420bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6421 SPIRVTypeInst ResType,
6422 MachineInstr &
I)
const {
6423 Register ImageReg =
I.getOperand(2).getReg();
6424 Register SamplerReg =
I.getOperand(3).getReg();
6425 Register CoordinateReg =
I.getOperand(4).getReg();
6426 ImageOperands ImOps;
6427 ImOps.Compare =
I.getOperand(5).getReg();
6428 if (
I.getNumOperands() > 6)
6429 ImOps.Offset =
I.getOperand(6).getReg();
6430 if (
I.getNumOperands() > 7)
6431 ImOps.MinLod =
I.getOperand(7).getReg();
6432 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6433 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6436bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6437 SPIRVTypeInst ResType,
6438 MachineInstr &
I)
const {
6439 Register ImageReg =
I.getOperand(2).getReg();
6440 Register CoordinateReg =
I.getOperand(3).getReg();
6441 Register LodReg =
I.getOperand(4).getReg();
6443 ImageOperands ImOps;
6445 if (
I.getNumOperands() > 5)
6446 ImOps.Offset =
I.getOperand(5).getReg();
6458 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6459 I.getDebugLoc(),
I, &ImOps);
6462bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6463 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6464 Register ImageReg =
I.getOperand(2).getReg();
6465 Register SamplerReg =
I.getOperand(3).getReg();
6466 Register CoordinateReg =
I.getOperand(4).getReg();
6467 ImageOperands ImOps;
6468 ImOps.Compare =
I.getOperand(5).getReg();
6469 if (
I.getNumOperands() > 6)
6470 ImOps.Offset =
I.getOperand(6).getReg();
6473 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6474 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6477bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6478 SPIRVTypeInst ResType,
6479 MachineInstr &
I)
const {
6480 Register ImageReg =
I.getOperand(2).getReg();
6481 Register SamplerReg =
I.getOperand(3).getReg();
6482 Register CoordinateReg =
I.getOperand(4).getReg();
6485 "ImageReg is not an image type.");
6490 ComponentOrCompareReg =
I.getOperand(5).getReg();
6491 OffsetReg =
I.getOperand(6).getReg();
6494 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6498 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6499 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6500 Dim != SPIRV::Dim::DIM_Rect) {
6502 "Gather operations are only supported for 2D, Cube, and Rect images.");
6509 if (!loadHandleBeforePosition(
6514 MachineIRBuilder MIRBuilder(
I);
6515 SPIRVTypeInst SampledImageType =
6520 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6528 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6530 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6532 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6537 .
addUse(ComponentOrCompareReg);
6539 uint32_t ImageOperands = 0;
6540 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6541 if (Dim == SPIRV::Dim::DIM_Cube) {
6543 "Gather operations with offset are not supported for Cube images.");
6547 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6549 ImageOperands |= SPIRV::ImageOperand::Offset;
6553 if (ImageOperands != 0) {
6554 MIB.
addImm(ImageOperands);
6556 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6564bool SPIRVInstructionSelector::generateImageReadOrFetch(
6567 const ImageOperands *ImOps)
const {
6570 "ImageReg is not an image type.");
6572 bool IsSignedInteger =
6577 bool IsFetch = (SampledOp.getImm() == 1);
6579 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6580 uint32_t ImageOperandsMask = 0;
6581 if (IsSignedInteger)
6582 ImageOperandsMask |= 0x1000;
6584 if (IsFetch && ImOps) {
6586 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6587 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6589 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6591 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6595 if (ImageOperandsMask != 0) {
6596 MIB.
addImm(ImageOperandsMask);
6597 if (IsFetch && ImOps) {
6600 if (ImOps->Offset &&
6601 (ImageOperandsMask &
6602 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6603 MIB.
addUse(*ImOps->Offset);
6612 SPIRVTypeInst SampledType =
6615 SPIRVTypeInst ReadType =
6616 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6617 bool ReadTypeMatchesResult = ReadType == ResType;
6619 Register ReadReg = ReadTypeMatchesResult
6625 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6631 BMI.constrainAllUses(
TII,
TRI, RBI);
6633 if (ReadTypeMatchesResult)
6646 if (ResultSize == 1) {
6655 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6658bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6659 SPIRVTypeInst ResType,
6660 MachineInstr &
I)
const {
6661 Register ResourcePtr =
I.getOperand(2).getReg();
6663 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6672 MachineIRBuilder MIRBuilder(
I);
6677 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6683 if (
I.getNumExplicitOperands() > 3) {
6684 Register IndexReg =
I.getOperand(3).getReg();
6691bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6692 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6697bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6698 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6699 Register ObjReg =
I.getOperand(2).getReg();
6700 if (!BuildCOPY(ResVReg, ObjReg,
I))
6710 decorateUsesAsNonUniform(ResVReg);
6714void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6717 {NonUniformReg,
nullptr}};
6718 llvm::SmallSet<Register, 8> Visited;
6719 while (WorkList.
size() > 0) {
6722 if (!Visited.
insert(CurrentReg).second)
6725 bool IsDecorated =
false;
6727 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6728 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6734 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6736 if (ResultReg == CurrentReg)
6744 MachineInstr &InsertPt =
6747 SPIRV::Decoration::NonUniformEXT, {});
6752bool SPIRVInstructionSelector::extractSubvector(
6754 MachineInstr &InsertionPoint)
const {
6756 [[maybe_unused]]
uint64_t InputSize =
6759 [[maybe_unused]]
bool IsLongVectorEXT =
6761 assert((InputSize > 1 || IsLongVectorEXT) &&
"The input must be a vector.");
6762 assert((ResultSize > 1 || IsLongVectorEXT) &&
"The result must be a vector.");
6763 assert(ResultSize < InputSize &&
6764 "Cannot extract more element than there are in the input.");
6771 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6780 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6782 TII.get(SPIRV::OpCompositeConstruct))
6786 for (
Register ComponentReg : ComponentRegisters)
6787 MIB.
addUse(ComponentReg);
6792bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6793 MachineInstr &
I)
const {
6800 Register ImageReg =
I.getOperand(1).getReg();
6808 Register CoordinateReg =
I.getOperand(2).getReg();
6809 Register DataReg =
I.getOperand(3).getReg();
6812 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6820Register SPIRVInstructionSelector::buildPointerToResource(
6821 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6822 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6823 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6825 if (ArraySize == 1) {
6826 SPIRVTypeInst PtrType =
6829 "SpirvResType did not have an explicit layout.");
6834 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6835 SPIRVTypeInst VarPointerType =
6838 VarPointerType, Set,
Binding, Name, MIRBuilder);
6840 SPIRVTypeInst ResPointerType =
6853bool SPIRVInstructionSelector::selectFirstBitSet16(
6854 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6855 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6857 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6861 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6864bool SPIRVInstructionSelector::selectFirstBitSet32(
6866 unsigned BitSetOpcode)
const {
6867 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6870 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6877bool SPIRVInstructionSelector::selectFirstBitSet64(
6879 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6893 if (ComponentCount > 2) {
6894 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6896 unsigned Opcode) ->
bool {
6897 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6901 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6905 MachineIRBuilder MIRBuilder(
I);
6907 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6911 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6917 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6927 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6928 SPIRV::OpVectorExtractDynamic))
6930 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6931 SPIRV::OpVectorExtractDynamic))
6935 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6936 TII.get(SPIRV::OpVectorShuffle))
6944 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6950 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6951 TII.get(SPIRV::OpVectorShuffle))
6959 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6979 SelectOp = SPIRV::OpSelectSISCond;
6980 AddOp = SPIRV::OpIAddS;
6988 SelectOp = SPIRV::OpSelectVIVCond;
6989 AddOp = SPIRV::OpIAddV;
6995 Register RegSecondaryOffset = Reg0;
6999 if (SwapPrimarySide) {
7000 PrimaryReg = LowReg;
7001 SecondaryReg = HighReg;
7002 RegPrimaryOffset = Reg0;
7003 RegSecondaryOffset = Reg32;
7008 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
7009 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
7014 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
7015 SPIRV::OpINotEqual))
7022 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
7023 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
7028 if (SwapPrimarySide) {
7030 if (!selectOpWithSrcs(RegAdd, ResType,
I,
7031 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
7042 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
7043 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
7048 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
7049 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
7052 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
7056bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
7057 SPIRVTypeInst ResType,
7059 bool IsSigned)
const {
7061 Register OpReg =
I.getOperand(2).getReg();
7064 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
7065 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
7069 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7071 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7073 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7076 return diagnoseUnsupported(
7078 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
7082bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
7083 SPIRVTypeInst ResType,
7084 MachineInstr &
I)
const {
7086 Register OpReg =
I.getOperand(2).getReg();
7091 unsigned ExtendOpcode = SPIRV::OpUConvert;
7092 unsigned BitSetOpcode = GL::FindILsb;
7096 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7098 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7100 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7103 return diagnoseUnsupported(
I,
7104 "spv_firstbitlow only supports 16,32,64 bits.");
7108bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
7109 SPIRVTypeInst ResType,
7110 MachineInstr &
I)
const {
7115 bool UseUntypedPointers =
7116 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7117 unsigned Opcode = UseUntypedPointers
7118 ? SPIRV::OpUntypedVariableLengthArrayINTEL
7119 : SPIRV::OpVariableLengthArrayINTEL;
7121 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
7127 if (UseUntypedPointers) {
7130 "untyped variable length array result must have a recorded element "
7135 MIB.
addUse(
I.getOperand(2).getReg());
7139 unsigned Alignment =
I.getOperand(3).getImm();
7153 while (!Worklist.
empty()) {
7155 switch (
T->getOpcode()) {
7156 case SPIRV::OpTypeInt:
7157 case SPIRV::OpTypeFloat:
7158 case SPIRV::OpTypePointer:
7160 case SPIRV::OpTypeVector:
7161 case SPIRV::OpTypeVectorIdEXT:
7162 case SPIRV::OpTypeMatrix:
7163 case SPIRV::OpTypeArray: {
7164 Register OperandReg =
T->getOperand(1).getReg();
7168 case SPIRV::OpTypeStruct:
7169 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7170 Register OperandReg =
T->getOperand(Idx).getReg();
7185 Register TypeReg = Ty->getOperand(0).getReg();
7186 if (!Visited.
insert(TypeReg).second)
7189 switch (Ty->getOpcode()) {
7190 case SPIRV::OpTypePointer:
7191 if (Ty->getOperand(1).getImm() == SPIRV::StorageClass::StorageBuffer)
7195 case SPIRV::OpTypeArray:
7196 case SPIRV::OpTypeRuntimeArray:
7199 case SPIRV::OpTypeStruct:
7200 for (
unsigned I = 1;
I < Ty->getNumOperands(); ++
I)
7216bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7217 assert(
I.getNumExplicitOperands() == 2);
7219 Register MsgReg =
I.getOperand(1).getReg();
7221 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7224 return diagnoseUnsupported(
7226 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7227 "scalar, pointer, vector, matrix, or aggregate of such types)");
7230 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7237bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7246 uint32_t MsgVal = ~0
u;
7247 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7248 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7251 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7254 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7261bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7262 SPIRVTypeInst ResType,
7263 MachineInstr &
I)
const {
7270 bool UseUntypedPointers =
7271 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7273 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7276 MachineIRBuilder MIRBuilder(
I);
7279 .
addImm(SPIRV::Extension::SPV_KHR_variable_pointers);
7281 .
addImm(SPIRV::Capability::VariablePointersStorageBuffer);
7284 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7287 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7291 if (UseUntypedPointers) {
7295 return diagnoseUnsupported(
7296 I,
"could not deduce the data type of an untyped variable");
7302 unsigned Alignment =
I.getOperand(2).getImm();
7309bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7314 const MachineInstr *PrevI =
I.getPrevNode();
7316 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7320 .
addMBB(
I.getOperand(0).getMBB())
7325 .
addMBB(
I.getOperand(0).getMBB())
7330bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7341 const MachineInstr *NextI =
I.getNextNode();
7343 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7349 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7351 .
addUse(
I.getOperand(0).getReg())
7352 .
addMBB(
I.getOperand(1).getMBB())
7358bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7359 MachineInstr &
I)
const {
7361 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7363 const unsigned NumOps =
I.getNumOperands();
7364 for (
unsigned i = 1; i <
NumOps; i += 2) {
7365 MIB.
addUse(
I.getOperand(i + 0).getReg());
7366 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7372bool SPIRVInstructionSelector::selectGlobalValue(
7373 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7375 MachineIRBuilder MIRBuilder(
I);
7376 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7379 std::string GlobalIdent;
7381 unsigned &
ID = UnnamedGlobalIDs[GV];
7383 ID = UnnamedGlobalIDs.
size();
7384 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7410 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7417 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7422 MachineInstrBuilder MIB1 =
7423 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7426 MachineInstrBuilder MIB2 =
7428 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7432 GR.
add(ConstVal, MIB2);
7440 MachineInstrBuilder MIB3 =
7441 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7444 GR.
add(ConstVal, MIB3);
7450 assert(NewReg != ResVReg);
7451 return BuildCOPY(ResVReg, NewReg,
I);
7461 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7464 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7470 SPIRVTypeInst ResType =
7474 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7479 if (
GlobalVar->isExternallyInitialized() &&
7480 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7481 constexpr unsigned ReadWriteINTEL = 3u;
7484 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7490bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7491 SPIRVTypeInst ResType,
7492 MachineInstr &
I)
const {
7494 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7502 MachineIRBuilder MIRBuilder(
I);
7507 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7510 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7512 .
add(
I.getOperand(1))
7526 APFloat::rmNearestTiesToEven, &LosesInfo);
7531 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
7541bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7542 SPIRVTypeInst ResType,
7543 MachineInstr &
I)
const {
7546 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7552 Register ExpReg =
I.getOperand(2).getReg();
7554 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7555 SPIRV::OpConvertSToF))
7557 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7564bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7565 SPIRVTypeInst ResType,
7566 MachineInstr &
I)
const {
7582 MachineIRBuilder MIRBuilder(
I);
7583 SPIRVTypeInst FloatType =
7587 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7600 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7601 const bool IsUntyped =
7602 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7604 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7605 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7606 : SPIRV::OpVariable))
7609 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7617 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7620 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7623 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7627 Register IntegralPartReg =
I.getOperand(1).getReg();
7630 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7640 assert(
false &&
"GLSL::Modf is deprecated.");
7651bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7652 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7653 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7654 MachineIRBuilder MIRBuilder(
I);
7655 const SPIRVTypeInst Vec3Ty =
7658 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7670 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7674 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7680 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7687 assert(
I.getOperand(2).isReg());
7688 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7692 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7703bool SPIRVInstructionSelector::loadBuiltinInputID(
7704 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7705 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7706 MachineIRBuilder MIRBuilder(
I);
7708 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7723 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7727 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7736SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7737 MachineInstr &
I)
const {
7738 MachineIRBuilder MIRBuilder(
I);
7749bool SPIRVInstructionSelector::loadHandleBeforePosition(
7750 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7751 MachineInstr &Pos)
const {
7754 Intrinsic::spv_resource_handlefrombinding);
7762 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7763 MachineIRBuilder MIRBuilder(HandleDef);
7764 SPIRVTypeInst VarType = ResType;
7765 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7767 if (IsStructuredBuffer) {
7776 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7779 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7780 ArraySize, IndexReg, Name, MIRBuilder);
7784 uint32_t LoadOpcode =
7785 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7795bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7796 MachineInstr &
I)
const {
7798 return diagnoseUnsupported(
7799 I,
"this instruction is only supported in shaders.");
7804InstructionSelector *
7808 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