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;
174 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
183 Register castPtrToPtrToInt(
Register Ptr, SPIRV::StorageClass::StorageClass SC,
187 bool selectAtomicPtrValue(
203 unsigned OpType)
const;
271 unsigned Opcode)
const;
275 unsigned Opcode)
const;
279 unsigned Opcode)
const;
283 unsigned Opcode)
const;
285 template <
bool Signed>
288 template <
bool Signed>
295 template <
typename PickOpcodeFn>
298 PickOpcodeFn &&PickOpcode)
const;
315 template <
typename PickOpcodeFn>
318 PickOpcodeFn &&PickOpcode)
const;
336 bool IsSigned)
const;
338 bool IsSigned,
unsigned Opcode)
const;
340 bool IsSigned)
const;
346 bool IsSigned)
const;
387 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
388 bool useMISrc =
true,
390 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
391 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
392 bool useMISrc =
true,
394 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
395 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
396 bool setMIFlags =
true,
bool useMISrc =
true,
398 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
399 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
400 bool useMISrc =
true,
403 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
404 MachineInstr &
I)
const;
406 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
407 MachineInstr &
I)
const;
409 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
412 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I,
unsigned Opcode)
const;
415 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
416 bool WithGroupSync)
const;
418 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
419 MachineInstr &
I)
const;
421 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
422 MachineInstr &
I)
const;
426 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
427 MachineInstr &
I)
const;
429 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
430 MachineInstr &
I)
const;
432 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
433 MachineInstr &
I)
const;
434 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
435 MachineInstr &
I)
const;
436 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
437 SPIRVTypeInst ResType,
438 MachineInstr &
I)
const;
439 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
440 MachineInstr &
I)
const;
443 std::optional<Register> LodReg = std::nullopt)
const;
444 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
445 MachineInstr &
I)
const;
446 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
447 MachineInstr &
I)
const;
448 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
449 MachineInstr &
I)
const;
450 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
451 MachineInstr &
I)
const;
452 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
453 MachineInstr &
I)
const;
454 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
455 MachineInstr &
I)
const;
456 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
457 MachineInstr &
I)
const;
458 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
459 SPIRVTypeInst ResType,
460 MachineInstr &
I)
const;
461 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
462 MachineInstr &
I)
const;
463 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
464 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
465 MachineInstr &
I)
const;
466 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
467 MachineInstr &
I)
const;
468 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
469 MachineInstr &
I)
const;
470 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
471 MachineInstr &
I)
const;
472 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
473 MachineInstr &
I)
const;
474 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
475 MachineInstr &
I)
const;
477 bool selectCopySign(
Register ResVReg, SPIRVTypeInst ResType,
478 MachineInstr &
I)
const;
480 bool selectLdexp(
Register ResVReg, SPIRVTypeInst ResType,
481 MachineInstr &
I)
const;
482 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
483 MachineInstr &
I)
const;
484 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
485 MachineInstr &
I)
const;
486 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
487 MachineInstr &
I,
const unsigned DPdOpCode)
const;
489 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
490 SPIRVTypeInst ResType =
nullptr)
const;
491 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
492 SPIRVTypeInst ResType =
nullptr)
const;
494 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
495 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
496 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
498 MachineInstr &
I)
const;
499 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
501 bool wrapIntoSpecConstantOp(MachineInstr &
I,
504 Register getUcharPtrTypeReg(MachineInstr &
I,
505 SPIRV::StorageClass::StorageClass SC)
const;
506 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
508 uint32_t Opcode)
const;
509 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
510 SPIRVTypeInst SrcPtrTy)
const;
511 Register buildPointerToResource(SPIRVTypeInst ResType,
512 SPIRV::StorageClass::StorageClass SC,
513 uint32_t Set, uint32_t
Binding,
514 uint32_t ArraySize,
Register IndexReg,
516 MachineIRBuilder MIRBuilder)
const;
517 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
518 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
519 Register &ReadReg, MachineInstr &InsertionPoint)
const;
520 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
523 const ImageOperands *ImOps =
nullptr)
const;
524 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
526 Register CoordinateReg,
const ImageOperands &ImOps,
529 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
530 Register ResVReg, SPIRVTypeInst ResType,
531 MachineInstr &
I)
const;
532 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
533 Register ResVReg, SPIRVTypeInst ResType,
534 MachineInstr &
I)
const;
535 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
536 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
537 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
538 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
541 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
542 Register SrcReg,
unsigned int Opcode,
543 std::function<
bool(
Register, SPIRVTypeInst,
544 MachineInstr &,
Register,
unsigned)>
548bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
550 if (
TET->getTargetExtName() ==
"spirv.Image") {
553 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
554 return TET->getTypeParameter(0)->isIntegerTy();
558#define GET_GLOBALISEL_IMPL
559#include "SPIRVGenGlobalISel.inc"
560#undef GET_GLOBALISEL_IMPL
566 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
569#include
"SPIRVGenGlobalISel.inc"
572#include
"SPIRVGenGlobalISel.inc"
584 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
589 if (HasVRegsReset == &MF)
604 for (
const auto &
MBB : MF) {
605 for (
const auto &
MI :
MBB) {
608 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
612 LLT DstType = MRI.
getType(DstReg);
614 LLT SrcType = MRI.
getType(SrcReg);
615 if (DstType != SrcType)
620 if (DstRC != SrcRC && SrcRC)
632 while (!Stack.empty()) {
637 switch (
MI->getOpcode()) {
638 case TargetOpcode::G_INTRINSIC:
639 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
640 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
643 if (IntrID != Intrinsic::spv_const_composite &&
644 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
648 case TargetOpcode::G_BUILD_VECTOR:
649 case TargetOpcode::G_SPLAT_VECTOR:
651 i < OpDef->getNumOperands(); i++) {
656 Stack.push_back(OpNestedDef);
659 case TargetOpcode::G_CONSTANT:
660 case TargetOpcode::G_FCONSTANT:
661 case TargetOpcode::G_IMPLICIT_DEF:
662 case SPIRV::OpConstantTrue:
663 case SPIRV::OpConstantFalse:
664 case SPIRV::OpConstantI:
665 case SPIRV::OpConstantF:
666 case SPIRV::OpConstantComposite:
667 case SPIRV::OpConstantCompositeContinuedINTEL:
668 case SPIRV::OpConstantSampler:
669 case SPIRV::OpConstantNull:
671 case SPIRV::OpPoisonKHR:
672 case SPIRV::OpConstantFunctionPointerINTEL:
699 case Intrinsic::spv_all:
700 case Intrinsic::spv_alloca:
701 case Intrinsic::spv_any:
702 case Intrinsic::spv_bitcast:
703 case Intrinsic::spv_const_composite:
704 case Intrinsic::spv_degrees:
705 case Intrinsic::spv_distance:
706 case Intrinsic::spv_extractelt:
707 case Intrinsic::spv_extractv:
708 case Intrinsic::spv_faceforward:
709 case Intrinsic::spv_fdot:
710 case Intrinsic::spv_firstbitlow:
711 case Intrinsic::spv_firstbitshigh:
712 case Intrinsic::spv_firstbituhigh:
713 case Intrinsic::spv_frac:
714 case Intrinsic::spv_gep:
715 case Intrinsic::spv_global_offset:
716 case Intrinsic::spv_global_size:
717 case Intrinsic::spv_group_id:
718 case Intrinsic::spv_insertelt:
719 case Intrinsic::spv_insertv:
720 case Intrinsic::spv_isinf:
721 case Intrinsic::spv_isnan:
722 case Intrinsic::spv_isfinite:
723 case Intrinsic::spv_isnormal:
724 case Intrinsic::spv_lerp:
725 case Intrinsic::spv_length:
726 case Intrinsic::spv_normalize:
727 case Intrinsic::spv_num_subgroups:
728 case Intrinsic::spv_num_workgroups:
729 case Intrinsic::spv_ptrcast:
730 case Intrinsic::spv_radians:
731 case Intrinsic::spv_reflect:
732 case Intrinsic::spv_refract:
733 case Intrinsic::spv_resource_getbasepointer:
734 case Intrinsic::spv_resource_getpointer:
735 case Intrinsic::spv_resource_handlefrombinding:
736 case Intrinsic::spv_resource_handlefromimplicitbinding:
737 case Intrinsic::spv_resource_nonuniformindex:
738 case Intrinsic::spv_resource_sample:
739 case Intrinsic::spv_rsqrt:
740 case Intrinsic::spv_saturate:
741 case Intrinsic::spv_sdot:
742 case Intrinsic::spv_sign:
743 case Intrinsic::spv_smoothstep:
744 case Intrinsic::spv_subgroup_id:
745 case Intrinsic::spv_subgroup_local_invocation_id:
746 case Intrinsic::spv_subgroup_max_size:
747 case Intrinsic::spv_subgroup_size:
748 case Intrinsic::spv_thread_id:
749 case Intrinsic::spv_thread_id_in_group:
750 case Intrinsic::spv_udot:
751 case Intrinsic::spv_undef:
752 case Intrinsic::spv_value_md:
753 case Intrinsic::spv_workgroup_size:
765 case SPIRV::OpTypeVoid:
766 case SPIRV::OpTypeBool:
767 case SPIRV::OpTypeInt:
768 case SPIRV::OpTypeFloat:
769 case SPIRV::OpTypeVector:
770 case SPIRV::OpTypeVectorIdEXT:
771 case SPIRV::OpTypeMatrix:
772 case SPIRV::OpTypeImage:
773 case SPIRV::OpTypeSampler:
774 case SPIRV::OpTypeSampledImage:
775 case SPIRV::OpTypeArray:
776 case SPIRV::OpTypeRuntimeArray:
777 case SPIRV::OpTypeStruct:
778 case SPIRV::OpTypeOpaque:
779 case SPIRV::OpTypePointer:
780 case SPIRV::OpTypeFunction:
781 case SPIRV::OpTypeEvent:
782 case SPIRV::OpTypeDeviceEvent:
783 case SPIRV::OpTypeReserveId:
784 case SPIRV::OpTypeQueue:
785 case SPIRV::OpTypePipe:
786 case SPIRV::OpTypeForwardPointer:
787 case SPIRV::OpTypePipeStorage:
788 case SPIRV::OpTypeNamedBarrier:
789 case SPIRV::OpTypeAccelerationStructureNV:
790 case SPIRV::OpTypeCooperativeMatrixNV:
791 case SPIRV::OpTypeCooperativeMatrixKHR:
801 if (
MI.getNumDefs() == 0)
804 for (
const auto &MO :
MI.all_defs()) {
806 if (
Reg.isPhysical()) {
811 if (
UseMI.getOpcode() != SPIRV::OpName) {
818 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
819 MI.isLifetimeMarker()) {
822 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
833 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
834 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
837 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
842 if (
MI.mayStore() ||
MI.isCall() ||
843 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
844 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
845 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
856 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
863void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
865 for (
const auto &MO :
MI.all_defs()) {
869 SmallVector<MachineInstr *, 4> UselessOpNames;
872 "There is still a use of the dead function.");
875 for (MachineInstr *OpNameMI : UselessOpNames) {
877 OpNameMI->eraseFromParent();
882void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
885 removeOpNamesForDeadMI(
MI);
886 MI.eraseFromParent();
889bool SPIRVInstructionSelector::select(MachineInstr &
I) {
890 resetVRegsType(*
I.getParent()->getParent());
892 assert(
I.getParent() &&
"Instruction should be in a basic block!");
893 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
898 removeDeadInstruction(
I);
905 if (Opcode == SPIRV::ASSIGN_TYPE) {
906 Register DstReg =
I.getOperand(0).getReg();
907 Register SrcReg =
I.getOperand(1).getReg();
910 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
911 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
912 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
913 Register SelectDstReg =
Def->getOperand(0).getReg();
914 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
916 assert(SuccessToSelectSelect);
918 Def->eraseFromParent();
925 bool Res = selectImpl(
I, *CoverageInfo);
927 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
928 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
932 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
944 }
else if (
I.getNumDefs() == 1) {
956 removeDeadInstruction(
I);
961 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
962 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
968 bool HasDefs =
I.getNumDefs() > 0;
971 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
972 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
973 if (spvSelect(ResVReg, ResType,
I)) {
975 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
986 case TargetOpcode::G_CONSTANT:
987 case TargetOpcode::G_FCONSTANT:
994 MachineInstr &
I)
const {
997 if (DstRC != SrcRC && SrcRC)
999 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
1006bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
1007 SPIRVTypeInst ResType,
1008 MachineInstr &
I)
const {
1009 const unsigned Opcode =
I.getOpcode();
1011 return selectImpl(
I, *CoverageInfo);
1013 case TargetOpcode::G_CONSTANT:
1014 case TargetOpcode::G_FCONSTANT:
1015 return selectConst(ResVReg, ResType,
I);
1016 case TargetOpcode::G_GLOBAL_VALUE:
1017 return selectGlobalValue(ResVReg,
I);
1018 case TargetOpcode::G_IMPLICIT_DEF:
1019 return selectOpUndef(ResVReg, ResType,
I);
1020 case TargetOpcode::G_FREEZE:
1021 return selectFreeze(ResVReg, ResType,
I);
1023 case TargetOpcode::G_INTRINSIC:
1024 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1025 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1026 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1027 return selectIntrinsic(ResVReg, ResType,
I);
1028 case TargetOpcode::G_BITREVERSE:
1029 return selectBitreverse(ResVReg, ResType,
I);
1031 case TargetOpcode::G_BUILD_VECTOR:
1032 return selectBuildVector(ResVReg, ResType,
I);
1033 case TargetOpcode::G_SPLAT_VECTOR:
1034 return selectSplatVector(ResVReg, ResType,
I);
1035 case TargetOpcode::G_CONCAT_VECTORS:
1036 return selectConcatVectors(ResVReg, ResType,
I);
1038 case TargetOpcode::G_SHUFFLE_VECTOR: {
1039 MachineBasicBlock &BB = *
I.getParent();
1040 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1043 .
addUse(
I.getOperand(1).getReg())
1044 .
addUse(
I.getOperand(2).getReg());
1045 for (
auto V :
I.getOperand(3).getShuffleMask())
1050 case TargetOpcode::G_MEMMOVE:
1051 case TargetOpcode::G_MEMCPY:
1052 case TargetOpcode::G_MEMCPY_INLINE:
1053 case TargetOpcode::G_MEMSET:
1054 case TargetOpcode::G_MEMSET_INLINE:
1055 return selectMemOperation(ResVReg,
I);
1057 case TargetOpcode::G_ICMP:
1058 return selectICmp(ResVReg, ResType,
I);
1059 case TargetOpcode::G_FCMP:
1060 return selectFCmp(ResVReg, ResType,
I);
1062 case TargetOpcode::G_FRAME_INDEX:
1063 return selectFrameIndex(ResVReg, ResType,
I);
1065 case TargetOpcode::G_LOAD:
1066 return selectLoad(ResVReg, ResType,
I);
1067 case TargetOpcode::G_STORE:
1068 return selectStore(
I);
1070 case TargetOpcode::G_BR:
1071 return selectBranch(
I);
1072 case TargetOpcode::G_BRCOND:
1073 return selectBranchCond(
I);
1075 case TargetOpcode::G_PHI:
1076 return selectPhi(ResVReg,
I);
1078 case TargetOpcode::G_FPTOSI:
1079 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1080 case TargetOpcode::G_FPTOUI:
1081 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1083 case TargetOpcode::G_FPTOSI_SAT:
1084 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1085 case TargetOpcode::G_FPTOUI_SAT:
1086 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1088 case TargetOpcode::G_SITOFP:
1089 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1090 case TargetOpcode::G_UITOFP:
1091 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1093 case TargetOpcode::G_CTPOP:
1094 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1095 case TargetOpcode::G_SMIN:
1096 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1097 case TargetOpcode::G_UMIN:
1098 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1100 case TargetOpcode::G_SMAX:
1101 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1102 case TargetOpcode::G_UMAX:
1103 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1105 case TargetOpcode::G_SCMP:
1106 return selectSUCmp(ResVReg, ResType,
I,
true);
1107 case TargetOpcode::G_UCMP:
1108 return selectSUCmp(ResVReg, ResType,
I,
false);
1109 case TargetOpcode::G_LROUND:
1110 case TargetOpcode::G_LLROUND: {
1113 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1115 regForLround, *(
I.getParent()->getParent()));
1117 CL::round, GL::Round,
false);
1119 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1126 case TargetOpcode::G_STRICT_FMA:
1127 case TargetOpcode::G_FMA: {
1130 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1133 .
addUse(
I.getOperand(1).getReg())
1134 .
addUse(
I.getOperand(2).getReg())
1135 .
addUse(
I.getOperand(3).getReg())
1140 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1143 case TargetOpcode::G_FLDEXP:
1144 case TargetOpcode::G_STRICT_FLDEXP:
1145 return selectLdexp(ResVReg, ResType,
I);
1147 case TargetOpcode::G_FPOW:
1148 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1149 case TargetOpcode::G_FPOWI:
1150 return selectFpowi(ResVReg, ResType,
I);
1152 case TargetOpcode::G_FEXP:
1153 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1154 case TargetOpcode::G_FEXP2:
1155 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1156 case TargetOpcode::G_FEXP10:
1157 return selectExp10(ResVReg, ResType,
I);
1159 case TargetOpcode::G_FMODF:
1160 return selectModf(ResVReg, ResType,
I);
1161 case TargetOpcode::G_FSINCOS:
1162 return selectSincos(ResVReg, ResType,
I);
1164 case TargetOpcode::G_FLOG:
1165 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1166 case TargetOpcode::G_FLOG2:
1167 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1168 case TargetOpcode::G_FLOG10:
1169 return selectLog10(ResVReg, ResType,
I);
1171 case TargetOpcode::G_FABS:
1172 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1173 case TargetOpcode::G_ABS:
1174 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1176 case TargetOpcode::G_FMINNUM:
1177 case TargetOpcode::G_FMINIMUM:
1178 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1179 case TargetOpcode::G_FMAXNUM:
1180 case TargetOpcode::G_FMAXIMUM:
1181 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1183 case TargetOpcode::G_FCOPYSIGN:
1184 return selectCopySign(ResVReg, ResType,
I);
1186 case TargetOpcode::G_FCEIL:
1187 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1188 case TargetOpcode::G_FFLOOR:
1189 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1191 case TargetOpcode::G_FCOS:
1192 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1193 case TargetOpcode::G_FSIN:
1194 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1195 case TargetOpcode::G_FTAN:
1196 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1197 case TargetOpcode::G_FACOS:
1198 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1199 case TargetOpcode::G_FASIN:
1200 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1201 case TargetOpcode::G_FATAN:
1202 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1203 case TargetOpcode::G_FATAN2:
1204 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1205 case TargetOpcode::G_FCOSH:
1206 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1207 case TargetOpcode::G_FSINH:
1208 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1209 case TargetOpcode::G_FTANH:
1210 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1212 case TargetOpcode::G_STRICT_FSQRT:
1213 case TargetOpcode::G_FSQRT:
1214 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1216 case TargetOpcode::G_CTTZ:
1217 case TargetOpcode::G_CTTZ_ZERO_POISON:
1218 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1219 case TargetOpcode::G_CTLZ:
1220 case TargetOpcode::G_CTLZ_ZERO_POISON:
1221 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1223 case TargetOpcode::G_INTRINSIC_ROUND:
1224 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1225 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1226 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1227 case TargetOpcode::G_INTRINSIC_TRUNC:
1228 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1229 case TargetOpcode::G_FRINT:
1230 case TargetOpcode::G_FNEARBYINT:
1231 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1233 case TargetOpcode::G_SMULH:
1234 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1235 case TargetOpcode::G_UMULH:
1236 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1238 case TargetOpcode::G_SADDSAT:
1239 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1240 case TargetOpcode::G_UADDSAT:
1241 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1242 case TargetOpcode::G_SSUBSAT:
1243 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1244 case TargetOpcode::G_USUBSAT:
1245 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1247 case TargetOpcode::G_FFREXP:
1248 return selectFrexp(ResVReg, ResType,
I);
1250 case TargetOpcode::G_UADDO:
1251 return selectOverflowArith(ResVReg, ResType,
I,
1253 : SPIRV::OpIAddCarryS);
1254 case TargetOpcode::G_USUBO:
1255 return selectOverflowArith(ResVReg, ResType,
I,
1257 : SPIRV::OpISubBorrowS);
1258 case TargetOpcode::G_UMULO:
1259 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1260 case TargetOpcode::G_SMULO:
1261 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1263 case TargetOpcode::G_SEXT:
1264 return selectExt(ResVReg, ResType,
I,
true);
1265 case TargetOpcode::G_ANYEXT:
1266 case TargetOpcode::G_ZEXT:
1267 return selectExt(ResVReg, ResType,
I,
false);
1268 case TargetOpcode::G_TRUNC:
1269 return selectTrunc(ResVReg, ResType,
I);
1270 case TargetOpcode::G_FPTRUNC:
1271 case TargetOpcode::G_FPEXT:
1272 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1274 case TargetOpcode::G_PTRTOINT:
1275 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1276 case TargetOpcode::G_INTTOPTR:
1277 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1278 case TargetOpcode::G_BITCAST:
1279 return selectBitcast(ResVReg, ResType,
I);
1280 case TargetOpcode::G_ADDRSPACE_CAST:
1281 return selectAddrSpaceCast(ResVReg, ResType,
I);
1282 case TargetOpcode::G_PTRMASK:
1283 return selectPtrMask(ResVReg, ResType,
I);
1284 case TargetOpcode::G_PTR_ADD: {
1286 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1290 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1291 (*II).getOpcode() == TargetOpcode::COPY ||
1292 (*II).getOpcode() == SPIRV::OpVariable ||
1293 (*II).getOpcode() == SPIRV::OpUntypedVariableKHR) &&
1294 getImm(
I.getOperand(2), MRI));
1296 bool IsGVInit =
false;
1300 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1301 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1302 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1303 (*UseIt).getOpcode() == SPIRV::OpVariable ||
1304 (*UseIt).getOpcode() == SPIRV::OpUntypedVariableKHR) {
1316 const bool UseUntypedPointers =
1317 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1318 if (UseUntypedPointers) {
1319 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1322 .
addImm(
static_cast<uint32_t
>(
1323 SPIRV::Opcode::UntypedInBoundsPtrAccessChainKHR))
1326 .
addUse(
I.getOperand(2).getReg())
1333 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1345 return diagnoseUnsupported(
1346 I,
"incompatible result and operand types in a bitcast");
1348 MachineInstrBuilder MIB =
1349 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1356 : SPIRV::OpInBoundsPtrAccessChain))
1360 .
addUse(
I.getOperand(2).getReg())
1363 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1367 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1369 .
addUse(
I.getOperand(2).getReg())
1378 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1381 .
addImm(
static_cast<uint32_t
>(
1382 SPIRV::Opcode::InBoundsPtrAccessChain))
1385 .
addUse(
I.getOperand(2).getReg());
1390 case TargetOpcode::G_ATOMICRMW_OR:
1391 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1392 case TargetOpcode::G_ATOMICRMW_ADD:
1393 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1394 case TargetOpcode::G_ATOMICRMW_AND:
1395 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1396 case TargetOpcode::G_ATOMICRMW_MAX:
1397 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1398 case TargetOpcode::G_ATOMICRMW_MIN:
1399 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1400 case TargetOpcode::G_ATOMICRMW_SUB:
1401 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1402 case TargetOpcode::G_ATOMICRMW_XOR:
1403 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1404 case TargetOpcode::G_ATOMICRMW_UMAX:
1405 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1406 case TargetOpcode::G_ATOMICRMW_UMIN:
1407 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1408 case TargetOpcode::G_ATOMICRMW_XCHG:
1409 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1411 case TargetOpcode::G_ATOMICRMW_FADD:
1412 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1413 case TargetOpcode::G_ATOMICRMW_FSUB:
1415 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1417 : SPIRV::OpFNegate);
1418 case TargetOpcode::G_ATOMICRMW_FMIN:
1419 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1420 case TargetOpcode::G_ATOMICRMW_FMAX:
1421 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1423 case TargetOpcode::G_FENCE:
1424 return selectFence(
I);
1426 case TargetOpcode::G_STACKSAVE:
1427 return selectStackSave(ResVReg, ResType,
I);
1428 case TargetOpcode::G_STACKRESTORE:
1429 return selectStackRestore(
I);
1431 case TargetOpcode::G_UNMERGE_VALUES:
1434 case TargetOpcode::G_TRAP:
1435 case TargetOpcode::G_UBSANTRAP:
1436 return selectTrap(
I);
1441 case TargetOpcode::DBG_LABEL:
1443 case TargetOpcode::G_DEBUGTRAP:
1444 return selectDebugTrap(ResVReg, ResType,
I);
1445 case TargetOpcode::G_PREFETCH:
1446 return selectPrefetch(
I);
1453bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1454 SPIRVTypeInst ResType,
1455 MachineInstr &
I)
const {
1456 unsigned Opcode = SPIRV::OpNop;
1463bool SPIRVInstructionSelector::selectPrefetch(MachineInstr &
I)
const {
1472 MachineIRBuilder MIRBuilder(
I);
1474 const SPIRVTypeInst PointerSizeType =
1482 Register AddrVal =
I.getOperand(0).getReg();
1485 return selectExtInst(ExtReg, GR.
getOpTypeVoid(MIRBuilder),
I, CL::prefetch,
1487 {AddrVal, ConstIntOne});
1492bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1493 SPIRVTypeInst ResType,
1495 GL::GLSLExtInst GLInst,
1496 bool setMIFlags,
bool useMISrc,
1499 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1500 return diagnoseUnsupported(
1502 "this instruction is only supported with the GLSL extended instruction "
1504 return selectExtInst(ResVReg, ResType,
I,
1505 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1506 setMIFlags, useMISrc, SrcRegs);
1509bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1510 SPIRVTypeInst ResType,
1512 CL::OpenCLExtInst CLInst,
1513 bool setMIFlags,
bool useMISrc,
1515 return selectExtInst(ResVReg, ResType,
I,
1516 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1517 setMIFlags, useMISrc, SrcRegs);
1520bool SPIRVInstructionSelector::selectExtInst(
1521 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1522 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1524 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1525 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1526 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1530bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1531 SPIRVTypeInst ResType,
1534 bool setMIFlags,
bool useMISrc,
1537 for (
const auto &[InstructionSet, Opcode] : Insts) {
1541 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1544 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1549 const unsigned NumOps =
I.getNumOperands();
1552 I.getOperand(Index).getType() ==
1553 MachineOperand::MachineOperandType::MO_IntrinsicID)
1556 MIB.
add(
I.getOperand(Index));
1568bool SPIRVInstructionSelector::selectCopySign(
Register ResVReg,
1569 SPIRVTypeInst ResType,
1570 MachineInstr &
I)
const {
1572 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1577 Register MagnitudeReg =
I.getOperand(1).getReg();
1578 Register SignReg =
I.getOperand(2).getReg();
1586 unsigned AndOpcode, OrOpcode;
1587 if (ComponentCount > 1) {
1591 AndOpcode = SPIRV::OpBitwiseAndV;
1592 OrOpcode = SPIRV::OpBitwiseOrV;
1596 AndOpcode = SPIRV::OpBitwiseAndS;
1597 OrOpcode = SPIRV::OpBitwiseOrS;
1603 return selectOpWithSrcs(ResReg, IntType,
I, SrcRegs, Opcode);
1606 Register MagnitudeInt, SignInt, MagnitudeBits, SignBits, CombinedInt;
1607 if (!EmitBitOp(MagnitudeInt, {MagnitudeReg}, SPIRV::OpBitcast) ||
1608 !EmitBitOp(SignInt, {SignReg}, SPIRV::OpBitcast) ||
1609 !EmitBitOp(MagnitudeBits, {MagnitudeInt, NotSignMask}, AndOpcode) ||
1610 !EmitBitOp(SignBits, {SignInt, SignMask}, AndOpcode) ||
1611 !EmitBitOp(CombinedInt, {MagnitudeBits, SignBits}, OrOpcode))
1614 return selectOpWithSrcs(ResVReg, ResType,
I, {CombinedInt}, SPIRV::OpBitcast);
1617bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1618 SPIRVTypeInst ResType,
1619 MachineInstr &
I)
const {
1620 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1621 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1622 for (
const auto &Ex : ExtInsts) {
1623 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1624 uint32_t Opcode = Ex.second;
1628 MachineIRBuilder MIRBuilder(
I);
1631 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1638 const bool IsUntyped =
1639 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1641 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1642 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1643 : SPIRV::OpVariable))
1646 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1652 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1655 .
addImm(
static_cast<uint32_t
>(Ex.first))
1657 .
add(
I.getOperand(2))
1661 Register ExpResReg =
I.getOperand(1).getReg();
1663 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1673bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1674 SPIRVTypeInst ResType,
1675 MachineInstr &
I)
const {
1676 Register XReg =
I.getOperand(1).getReg();
1677 Register ExpReg =
I.getOperand(2).getReg();
1683 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1684 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1686 SPIRVTypeInst ExpVecType =
1690 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1691 TII.get(SPIRV::OpCompositeConstruct))
1694 for (
unsigned J = 0; J < NumElts; ++J)
1700 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1701 true,
false, {XReg, ExpReg});
1704bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1705 SPIRVTypeInst ResType,
1706 MachineInstr &
I)
const {
1707 Register CosResVReg =
I.getOperand(1).getReg();
1708 unsigned SrcIdx =
I.getNumExplicitDefs();
1713 MachineIRBuilder MIRBuilder(
I);
1715 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1722 const bool IsUntyped =
1723 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1725 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1726 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1727 : SPIRV::OpVariable))
1730 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1734 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1737 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1739 .
add(
I.getOperand(SrcIdx))
1743 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1751 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1754 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1756 .
add(
I.getOperand(SrcIdx))
1758 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1761 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1763 .
add(
I.getOperand(SrcIdx))
1770bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1771 SPIRVTypeInst ResType,
1774 unsigned Opcode)
const {
1775 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1785bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1786 SPIRVTypeInst ResType,
1789 unsigned Opcode)
const {
1790 MachineIRBuilder MIRBuilder(
I);
1792 Register OpReg =
I.getOperand(1).getReg();
1798 SPIRVTypeInst ExtType =
1805 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1809 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1812 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1815bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1816 SPIRVTypeInst ResType,
1819 unsigned Opcode)
const {
1820 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1823bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1824 SPIRVTypeInst ResType,
1827 unsigned Opcode)
const {
1828 MachineIRBuilder MIRBuilder(
I);
1837 SPIRVTypeInst WorkingType =
1844 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {SrcReg}, SPIRV::OpUConvert))
1848 if (!selectOpWithSrcs(LowCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1856 IsScalar ? SPIRV::OpShiftRightLogicalS : SPIRV::OpShiftRightLogicalV;
1858 if (!selectOpWithSrcs(Shift, SrcType,
I, {SrcReg, ShiftAmount}, ShiftOp))
1862 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {Shift}, SPIRV::OpUConvert))
1866 if (!selectOpWithSrcs(HighCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1871 if (!selectOpWithSrcs(Sum, WorkingType,
I, {HighCount, LowCount},
1872 IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV))
1876 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1877 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1880bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1881 SPIRVTypeInst ResType,
1883 unsigned Opcode)
const {
1888 if (!STI.getTargetTriple().isVulkanOS())
1889 return selectUnOp(ResVReg, ResType,
I, Opcode);
1891 Register OpReg =
I.getOperand(1).getReg();
1894 : SPIRV::OpUConvert;
1898 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1900 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1902 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1904 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1908bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1909 SPIRVTypeInst ResType,
1911 unsigned Opcode)
const {
1913 Register SrcReg =
I.getOperand(1).getReg();
1918 unsigned DefOpCode = DefIt->getOpcode();
1919 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1922 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1923 DefOpCode = VRD->getOpcode();
1925 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1926 DefOpCode == TargetOpcode::G_CONSTANT ||
1927 DefOpCode == SPIRV::OpVariable ||
1928 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1929 DefOpCode == SPIRV::OpConstantI) {
1935 uint32_t SpecOpcode = 0;
1937 case SPIRV::OpConvertPtrToU:
1938 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1940 case SPIRV::OpConvertUToPtr:
1941 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1946 TII.get(SPIRV::OpSpecConstantOp))
1956 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1960bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1961 SPIRVTypeInst ResType,
1962 MachineInstr &
I)
const {
1963 Register OpReg =
I.getOperand(1).getReg();
1964 SPIRVTypeInst OpType =
1967 return diagnoseUnsupported(
1968 I,
"incompatible result and operand types in a bitcast");
1969 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1980 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1981 if (
MemOp->isNonTemporal())
1982 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1984 if (!ST->isShader() &&
MemOp->getAlign().value())
1985 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1989 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1990 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1994 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1996 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
2000 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
2004 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
2006 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
2018 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2020 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2022 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2026bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2027 SPIRVTypeInst ResType,
2028 MachineInstr &
I)
const {
2030 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2035 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2036 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2038 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2040 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2044 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2048 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2049 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2050 I.getDebugLoc(),
I);
2054 MachineIRBuilder MIRBuilder(
I);
2056 if (
I.getNumMemOperands()) {
2057 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2058 if (MemOp->isAtomic())
2059 return selectAtomicLoad(ResVReg, ResType,
I);
2062 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2066 if (!
I.getNumMemOperands()) {
2067 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2069 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2078Register SPIRVInstructionSelector::createPtrSizedIntReg(
2079 MachineIRBuilder &MIRBuilder)
const {
2080 SPIRVTypeInst IntType =
2090SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2091 MachineIRBuilder &MIRBuilder)
const {
2092 SPIRVTypeInst IntType =
2094 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2095 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2103Register SPIRVInstructionSelector::castPtrToPtrToInt(
2104 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2105 MachineIRBuilder &MIRBuilder)
const {
2106 SPIRVTypeInst IntType =
2108 SPIRVTypeInst PtrType =
2122bool SPIRVInstructionSelector::selectAtomicPtrValue(
2123 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2124 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2132 Register IntResult = EmitAtomic(IntType);
2134 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2142bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2143 SPIRVTypeInst ResType,
2144 MachineInstr &
I)
const {
2145 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2148 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2151 return diagnoseUnsupported(
2152 I,
"Lowering to SPIR-V of atomic load is only "
2153 "allowed for integer, floating point or pointer types");
2155 assert(
I.getNumMemOperands());
2156 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2157 assert(MemOp.isAtomic());
2159 uint32_t
Scope =
static_cast<uint32_t
>(
2161 Register ScopeReg = buildI32Constant(Scope,
I);
2167 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2168 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2171 Register MemSemReg = buildI32Constant(Sem,
I);
2173 MachineIRBuilder MIRBuilder(
I);
2177 return diagnoseUnsupported(
2178 I,
"Lowering to SPIR-V of atomic load is only "
2179 "allowed for pointer types for physical addressing model");
2184 SPIRV::StorageClass::StorageClass SC =
2186 return selectAtomicPtrValue(
2187 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2188 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2189 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2200 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2211bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2213 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2214 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2219 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2220 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2222 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2227 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2231 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2232 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2233 SPIRVTypeInst SampledType =
2235 SPIRVTypeInst StoreValCompType =
2237 if (StoreValCompType && StoreValCompType != SampledType) {
2240 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2243 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2248 StoreVal = PackedReg;
2251 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2252 TII.get(SPIRV::OpImageWrite))
2258 if (sampledTypeIsSignedInteger(LLVMHandleType))
2261 BMI.constrainAllUses(
TII,
TRI, RBI);
2268 if (PointeeTy && PointeeTy->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
2269 StoreTy->
getOpcode() != SPIRV::OpTypeVectorIdEXT &&
2271 MachineInstr *StoreValDef =
getVRegDef(*MRI, StoreVal);
2283 if (
I.getNumMemOperands()) {
2284 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2285 if (MemOp->isAtomic())
2286 return selectAtomicStore(
I);
2293 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2294 PtrSC == SPIRV::StorageClass::Input ||
2295 PtrSC == SPIRV::StorageClass::PushConstant)
2296 return diagnoseUnsupported(
2297 I,
"store into a read-only SPIR-V storage class is not allowed");
2299 MachineIRBuilder MIRBuilder(
I);
2301 if (!
I.getNumMemOperands()) {
2302 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2304 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2313bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2314 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2317 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2318 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2323 if (!PointeeType && PtrType &&
2324 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2327 return diagnoseUnsupported(
I,
2328 "Lowering to SPIR-V of atomic store is only "
2329 "allowed for integer or floating point types");
2331 assert(
I.getNumMemOperands());
2332 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2333 assert(MemOp.isAtomic());
2335 uint32_t
Scope =
static_cast<uint32_t
>(
2337 Register ScopeReg = buildI32Constant(Scope,
I);
2343 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2344 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2347 Register MemSemReg = buildI32Constant(Sem,
I);
2348 MachineIRBuilder MIRBuilder(
I);
2352 return diagnoseUnsupported(
2353 I,
"Lowering to SPIR-V of atomic store is only "
2354 "allowed for pointer types for physical addressing model");
2359 SPIRV::StorageClass::StorageClass SC =
2361 return selectAtomicPtrValue(
2362 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2364 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2377 return diagnoseUnsupported(
I,
2378 "Lowering to SPIR-V of atomic store is only "
2379 "allowed for integer or floating point types");
2381 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2391bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2392 SPIRVTypeInst ResType,
2393 MachineInstr &
I)
const {
2394 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2402 const Register PtrsReg =
I.getOperand(2).getReg();
2403 const uint32_t
Alignment =
I.getOperand(3).getImm();
2404 const Register MaskReg =
I.getOperand(4).getReg();
2405 const Register PassthruReg =
I.getOperand(5).getReg();
2406 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2410 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2421bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2422 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2429 const Register ValuesReg =
I.getOperand(1).getReg();
2430 const Register PtrsReg =
I.getOperand(2).getReg();
2431 const uint32_t
Alignment =
I.getOperand(3).getImm();
2432 const Register MaskReg =
I.getOperand(4).getReg();
2433 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2437 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2446bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2447 const Twine &
Msg)
const {
2448 const Function &
F =
I.getMF()->getFunction();
2449 F.getContext().diagnose(
2450 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2454bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2455 SPIRVTypeInst ResType,
2456 MachineInstr &
I)
const {
2457 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2458 return diagnoseUnsupported(
2459 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2460 "SPIR-V extension: SPV_INTEL_variable_length_array");
2462 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2469bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2470 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2471 return diagnoseUnsupported(
2473 "llvm.stackrestore intrinsic: this instruction requires the following "
2474 "SPIR-V extension: SPV_INTEL_variable_length_array");
2475 if (!
I.getOperand(0).isReg())
2478 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2479 .
addUse(
I.getOperand(0).getReg())
2485SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2486 MachineIRBuilder MIRBuilder(
I);
2487 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2494 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2498 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2499 Type *ArrTy = ArrayType::get(ValTy, Num);
2501 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2504 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2515 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2516 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2517 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2518 : SPIRV::OpVariable))
2521 .
addImm(SPIRV::StorageClass::UniformConstant);
2534bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2537 Register DstReg =
I.getOperand(0).getReg();
2541 return diagnoseUnsupported(
2542 I,
"OpCopyMemory requires operands to have the same type");
2547 return diagnoseUnsupported(
2548 I,
"Unable to determine pointee type size for OpCopyMemory");
2549 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2550 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2551 return diagnoseUnsupported(
2552 I,
"OpCopyMemory requires the size to match the pointee type size");
2553 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2556 if (
I.getNumMemOperands()) {
2557 MachineIRBuilder MIRBuilder(
I);
2564bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2567 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2568 .
addUse(
I.getOperand(0).getReg())
2570 .
addUse(
I.getOperand(2).getReg());
2571 if (
I.getNumMemOperands()) {
2572 MachineIRBuilder MIRBuilder(
I);
2579bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2580 MachineInstr &
I)
const {
2582 Register SizeReg =
I.getOperand(2).getReg();
2584 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2588 Register SrcReg =
I.getOperand(1).getReg();
2589 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2590 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2591 Register VarReg = getOrCreateMemSetGlobal(
I);
2594 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2596 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2598 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2602 if (!selectCopyMemory(
I, SrcReg))
2605 if (!selectCopyMemorySized(
I, SrcReg))
2608 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2609 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2614bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2615 SPIRVTypeInst ResType,
2618 unsigned NegateOpcode)
const {
2620 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2621 uint32_t
Scope =
static_cast<uint32_t
>(
2623 MemOp->getSyncScopeID()));
2624 Register ScopeReg = buildI32Constant(Scope,
I);
2626 Register Ptr =
I.getOperand(1).getReg();
2627 uint32_t ScSem =
static_cast<uint32_t
>(
2631 Register MemSemReg = buildI32Constant(
2635 Register ValueReg =
I.getOperand(2).getReg();
2636 if (NegateOpcode != 0) {
2639 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2645 if (NewOpcode != SPIRV::OpAtomicExchange)
2646 return diagnoseUnsupported(
2647 I,
"Lowering to SPIR-V of this atomic operation is not "
2648 "allowed for pointer types");
2650 return diagnoseUnsupported(
2651 I,
"Lowering to SPIR-V of atomic exchange is only "
2652 "allowed for pointer types for physical addressing model");
2659 MachineIRBuilder MIRBuilder(
I);
2661 return selectAtomicPtrValue(
2662 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2664 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2665 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2666 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2674 return ExchangeResReg;
2678 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2689bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2690 unsigned ArgI =
I.getNumOperands() - 1;
2692 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2693 SPIRVTypeInst SrcType =
2697 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2701 unsigned CurrentIndex = 0;
2702 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2703 Register ResVReg =
I.getOperand(i).getReg();
2706 LLT ResLLT = MRI->
getType(ResVReg);
2712 ResType = ScalarType;
2721 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2727 for (
unsigned j = 0;
j < NumElements; ++
j) {
2728 MIB.
addImm(CurrentIndex + j);
2730 CurrentIndex += NumElements;
2734 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2746bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2749 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
2753 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2755 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2762bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2763 SPIRVTypeInst ResType,
2765 unsigned Opcode)
const {
2766 Type *ResTy =
nullptr;
2769 return diagnoseUnsupported(
2771 "Not enough info to select the arithmetic with overflow instruction");
2773 return diagnoseUnsupported(
I,
2774 "Expect struct type result for the arithmetic "
2775 "with overflow instruction");
2781 MachineIRBuilder MIRBuilder(
I);
2783 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2784 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2791 Register ZeroReg = buildZerosVal(ResType,
I);
2796 if (ResName.
size() > 0)
2804 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2805 MIB.
addUse(
I.getOperand(i).getReg());
2810 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2811 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2813 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2814 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2821 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2822 .
addDef(
I.getOperand(1).getReg())
2830bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2831 SPIRVTypeInst ResType,
2832 MachineInstr &
I)
const {
2834 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2835 Register Ptr =
I.getOperand(2).getReg();
2836 Register ScopeReg =
I.getOperand(5).getReg();
2837 Register MemSemEqReg =
I.getOperand(6).getReg();
2838 Register MemSemNeqReg =
I.getOperand(7).getReg();
2840 Register Val =
I.getOperand(4).getReg();
2844 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2863 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2870 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2882 case SPIRV::StorageClass::DeviceOnlyINTEL:
2883 case SPIRV::StorageClass::HostOnlyINTEL:
2892 bool IsGRef =
false;
2893 bool IsAllowedRefs =
2895 unsigned Opcode = It.getOpcode();
2896 if (Opcode == SPIRV::OpConstantComposite ||
2897 Opcode == SPIRV::OpSpecConstantComposite ||
2898 Opcode == SPIRV::OpVariable ||
2899 Opcode == SPIRV::OpUntypedVariableKHR ||
2900 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2901 return IsGRef = true;
2902 return Opcode == SPIRV::OpName;
2904 return IsAllowedRefs && IsGRef;
2907Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2908 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2910 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2914SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2916 uint32_t Opcode)
const {
2917 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2918 TII.get(SPIRV::OpSpecConstantOp))
2926SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2927 SPIRVTypeInst SrcPtrTy)
const {
2928 SPIRVTypeInst GenericPtrTy =
2932 SPIRV::StorageClass::Generic),
2936 MachineInstrBuilder MIB = buildSpecConstantOp(
2938 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2948bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2949 SPIRVTypeInst ResType,
2950 MachineInstr &
I)
const {
2954 Register SrcPtr =
I.getOperand(1).getReg();
2959 return BuildCOPY(ResVReg, SrcPtr,
I);
2969 unsigned SpecOpcode = [&]() ->
unsigned {
2970 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2971 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2973 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2975 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
2983 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2985 .constrainAllUses(
TII,
TRI, RBI);
2987 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2989 buildSpecConstantOp(
2991 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2992 .constrainAllUses(
TII,
TRI, RBI);
2999 return BuildCOPY(ResVReg, SrcPtr,
I);
3001 if ((SrcSC == SPIRV::StorageClass::Function &&
3002 DstSC == SPIRV::StorageClass::Private) ||
3003 (DstSC == SPIRV::StorageClass::Function &&
3004 SrcSC == SPIRV::StorageClass::Private))
3005 return BuildCOPY(ResVReg, SrcPtr,
I);
3009 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3012 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3015 SPIRVTypeInst GenericPtrTy =
3034 return selectUnOp(ResVReg, ResType,
I,
3035 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
3037 return selectUnOp(ResVReg, ResType,
I,
3038 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
3040 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3042 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3052bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3053 SPIRVTypeInst ResType,
3054 MachineInstr &
I)
const {
3056 return diagnoseUnsupported(
3057 I,
"G_PTRMASK is not supported with logical SPIR-V");
3062 Register PtrReg =
I.getOperand(1).getReg();
3063 Register MaskReg =
I.getOperand(2).getReg();
3082 ? SPIRV::OpBitwiseAndV
3083 : SPIRV::OpBitwiseAndS;
3106 return SPIRV::OpFOrdEqual;
3108 return SPIRV::OpFOrdGreaterThanEqual;
3110 return SPIRV::OpFOrdGreaterThan;
3112 return SPIRV::OpFOrdLessThanEqual;
3114 return SPIRV::OpFOrdLessThan;
3116 return SPIRV::OpFOrdNotEqual;
3118 return SPIRV::OpOrdered;
3120 return SPIRV::OpFUnordEqual;
3122 return SPIRV::OpFUnordGreaterThanEqual;
3124 return SPIRV::OpFUnordGreaterThan;
3126 return SPIRV::OpFUnordLessThanEqual;
3128 return SPIRV::OpFUnordLessThan;
3130 return SPIRV::OpFUnordNotEqual;
3132 return SPIRV::OpUnordered;
3142 return SPIRV::OpIEqual;
3144 return SPIRV::OpINotEqual;
3146 return SPIRV::OpSGreaterThanEqual;
3148 return SPIRV::OpSGreaterThan;
3150 return SPIRV::OpSLessThanEqual;
3152 return SPIRV::OpSLessThan;
3154 return SPIRV::OpUGreaterThanEqual;
3156 return SPIRV::OpUGreaterThan;
3158 return SPIRV::OpULessThanEqual;
3160 return SPIRV::OpULessThan;
3169 return SPIRV::OpPtrEqual;
3171 return SPIRV::OpPtrNotEqual;
3182 return SPIRV::OpLogicalEqual;
3184 return SPIRV::OpLogicalNotEqual;
3222bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3223 SPIRVTypeInst ResType,
3225 unsigned OpAnyOrAll)
const {
3226 assert(
I.getNumOperands() == 3);
3227 assert(
I.getOperand(2).isReg());
3229 Register InputRegister =
I.getOperand(2).getReg();
3232 assert(InputType &&
"VReg has no type assigned");
3236 assert(ResVReg ==
I.getOperand(0).getReg());
3237 return BuildCOPY(ResVReg, InputRegister,
I);
3241 unsigned SpirvNotEqualId =
3242 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3244 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3249 IsBoolTy ? InputRegister
3257 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3259 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3276bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3277 SPIRVTypeInst ResType,
3278 MachineInstr &
I)
const {
3279 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3282bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3283 SPIRVTypeInst ResType,
3284 MachineInstr &
I)
const {
3285 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3289bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3290 SPIRVTypeInst ResType,
3291 MachineInstr &
I)
const {
3292 assert(
I.getNumOperands() == 4);
3293 assert(
I.getOperand(2).isReg());
3294 assert(
I.getOperand(3).isReg());
3296 [[maybe_unused]] SPIRVTypeInst VecType =
3301 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3302 "dot product requires either a vector of at least 2 components or"
3303 " the SPV_EXT_long vector extension.");
3305 [[maybe_unused]] SPIRVTypeInst EltType =
3314 .
addUse(
I.getOperand(2).getReg())
3315 .
addUse(
I.getOperand(3).getReg())
3320bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3321 SPIRVTypeInst ResType,
3324 assert(
I.getNumOperands() == 4);
3325 assert(
I.getOperand(2).isReg());
3326 assert(
I.getOperand(3).isReg());
3329 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3333 .
addUse(
I.getOperand(2).getReg())
3334 .
addUse(
I.getOperand(3).getReg())
3341bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3342 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3343 assert(
I.getNumOperands() == 4);
3344 assert(
I.getOperand(2).isReg());
3345 assert(
I.getOperand(3).isReg());
3349 Register Vec0 =
I.getOperand(2).getReg();
3350 Register Vec1 =
I.getOperand(3).getReg();
3354 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3363 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3364 "dot product requires either a vector of at least 2 components "
3365 "or the SPV_EXT_long_vector extension.");
3368 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3378 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3389 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3401bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3402 SPIRVTypeInst ResType,
3403 MachineInstr &
I)
const {
3405 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3408 .
addUse(
I.getOperand(2).getReg())
3413bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3414 SPIRVTypeInst ResType,
3415 MachineInstr &
I)
const {
3417 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3420 .
addUse(
I.getOperand(2).getReg())
3425bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3426 SPIRVTypeInst ResType,
3427 MachineInstr &
I)
const {
3429 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3432 .
addUse(
I.getOperand(2).getReg())
3437bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3438 SPIRVTypeInst ResType,
3439 MachineInstr &
I)
const {
3441 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3444 .
addUse(
I.getOperand(2).getReg())
3449template <
bool Signed>
3450bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3451 SPIRVTypeInst ResType,
3452 MachineInstr &
I)
const {
3453 assert(
I.getNumOperands() == 5);
3454 assert(
I.getOperand(2).isReg());
3455 assert(
I.getOperand(3).isReg());
3456 assert(
I.getOperand(4).isReg());
3459 Register Acc =
I.getOperand(2).getReg();
3463 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3465 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3470 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3473 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3485template <
bool Signed>
3486bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3487 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3488 assert(
I.getNumOperands() == 5);
3489 assert(
I.getOperand(2).isReg());
3490 assert(
I.getOperand(3).isReg());
3491 assert(
I.getOperand(4).isReg());
3494 Register Acc =
I.getOperand(2).getReg();
3500 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3504 for (
unsigned i = 0; i < 4; i++) {
3527 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3547 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3562bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3563 SPIRVTypeInst ResType,
3564 MachineInstr &
I)
const {
3565 assert(
I.getNumOperands() == 3);
3566 assert(
I.getOperand(2).isReg());
3568 Register VZero = buildZerosValF(ResType,
I);
3569 Register VOne = buildOnesValF(ResType,
I);
3571 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3574 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3576 .
addUse(
I.getOperand(2).getReg())
3583bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3584 SPIRVTypeInst ResType,
3585 MachineInstr &
I)
const {
3586 assert(
I.getNumOperands() == 3);
3587 assert(
I.getOperand(2).isReg());
3589 Register InputRegister =
I.getOperand(2).getReg();
3591 auto &
DL =
I.getDebugLoc();
3594 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3601 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3603 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3611 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3616 if (NeedsConversion) {
3617 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3628bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3629 SPIRVTypeInst ResType,
3631 unsigned Opcode)
const {
3635 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3641 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3642 BMI.addUse(
I.getOperand(J).getReg());
3649bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3652 bool WithGroupSync)
const {
3654 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3656 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3658 assert(((Scope != SPIRV::Scope::Workgroup) ||
3659 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3660 "Workgroup Scope must set WorkGroupMemory semantic "
3661 "in Barrier instruction");
3663 assert(((Scope != SPIRV::Scope::Device) ||
3664 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3665 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3666 "Device Scope must set UniformMemory and ImageMemory semantic "
3667 "in Barrier instruction");
3673 if (WithGroupSync) {
3674 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3678 Register ScopeReg = buildI32Constant(Scope,
I);
3679 Register MemSemReg = buildI32Constant(MemSem,
I);
3681 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3685bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3686 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3691 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3692 SPIRV::OpGroupNonUniformBallot))
3697 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3702 .
addImm(SPIRV::GroupOperation::Reduce)
3709bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3710 SPIRVTypeInst ResType,
3711 MachineInstr &
I)
const {
3716 Register InputReg =
I.getOperand(2).getReg();
3721 bool IsVector = NumElems > 1 ||
3722 (InputType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
3736 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3737 SPIRV::OpGroupNonUniformAllEqual);
3742 ElementResults.
reserve(NumElems);
3744 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3757 ElemInput = Extracted;
3763 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3774 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3785bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3786 SPIRVTypeInst ResType,
3787 MachineInstr &
I)
const {
3789 assert(
I.getNumOperands() == 3);
3791 auto Op =
I.getOperand(2);
3801 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3803 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3804 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3825 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3829 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3836bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3837 SPIRVTypeInst ResType,
3839 bool IsUnsigned)
const {
3840 return selectWaveReduce(
3841 ResVReg, ResType,
I, IsUnsigned,
3842 [&](
Register InputRegister,
bool IsUnsigned) {
3843 const bool IsFloatTy =
3845 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3846 : SPIRV::OpGroupNonUniformSMax;
3847 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3851bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3852 SPIRVTypeInst ResType,
3854 bool IsUnsigned)
const {
3855 return selectWaveReduce(
3856 ResVReg, ResType,
I, IsUnsigned,
3857 [&](
Register InputRegister,
bool IsUnsigned) {
3858 const bool IsFloatTy =
3860 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3861 : SPIRV::OpGroupNonUniformSMin;
3862 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3866bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3867 SPIRVTypeInst ResType,
3868 MachineInstr &
I)
const {
3869 return selectWaveReduce(ResVReg, ResType,
I,
false,
3870 [&](
Register InputRegister,
bool IsUnsigned) {
3872 InputRegister, SPIRV::OpTypeFloat);
3873 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3874 : SPIRV::OpGroupNonUniformIAdd;
3878bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3879 SPIRVTypeInst ResType,
3880 MachineInstr &
I)
const {
3881 return selectWaveReduce(ResVReg, ResType,
I,
false,
3882 [&](
Register InputRegister,
bool IsUnsigned) {
3884 InputRegister, SPIRV::OpTypeFloat);
3885 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3886 : SPIRV::OpGroupNonUniformIMul;
3890template <
typename PickOpcodeFn>
3891bool SPIRVInstructionSelector::selectWaveReduce(
3892 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3893 PickOpcodeFn &&PickOpcode)
const {
3894 assert(
I.getNumOperands() == 3);
3895 assert(
I.getOperand(2).isReg());
3897 Register InputRegister =
I.getOperand(2).getReg();
3901 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3904 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3910 .
addImm(SPIRV::GroupOperation::Reduce)
3911 .
addUse(
I.getOperand(2).getReg())
3916bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3917 SPIRVTypeInst ResType,
3919 unsigned Opcode)
const {
3920 return selectWaveReduce(
3921 ResVReg, ResType,
I,
false,
3922 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3925bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3926 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3927 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3928 [&](
Register InputRegister,
bool IsUnsigned) {
3930 InputRegister, SPIRV::OpTypeFloat);
3932 ? SPIRV::OpGroupNonUniformFAdd
3933 : SPIRV::OpGroupNonUniformIAdd;
3937bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3938 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3939 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3940 [&](
Register InputRegister,
bool IsUnsigned) {
3942 InputRegister, SPIRV::OpTypeFloat);
3944 ? SPIRV::OpGroupNonUniformFMul
3945 : SPIRV::OpGroupNonUniformIMul;
3949template <
typename PickOpcodeFn>
3950bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3951 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3952 PickOpcodeFn &&PickOpcode)
const {
3953 assert(
I.getNumOperands() == 3);
3954 assert(
I.getOperand(2).isReg());
3956 Register InputRegister =
I.getOperand(2).getReg();
3960 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3963 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3969 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3970 .
addUse(
I.getOperand(2).getReg())
3975bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3976 SPIRVTypeInst ResType,
3979 assert(
I.getNumOperands() == 3);
3980 assert(
I.getOperand(2).isReg());
3982 Register InputRegister =
I.getOperand(2).getReg();
3988 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3999bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
4000 SPIRVTypeInst ResType,
4007 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
4012 : SPIRV::OpUConvert;
4014 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4017 ShiftOp = SPIRV::OpShiftRightLogicalV;
4022 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4023 TII.get(SPIRV::OpConstantComposite))
4026 for (
unsigned It = 0; It <
N; ++It)
4030 ShiftConst = CompositeReg;
4035 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
4040 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
4045 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
4050 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
4053bool SPIRVInstructionSelector::handle64BitOverflow(
4055 unsigned int Opcode,
4062 "handle64BitOverflow should only be used for integer types");
4064 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4066 MachineIRBuilder MIRBuilder(
I);
4068 SPIRVTypeInst I64x2Type =
4070 SPIRVTypeInst Vec2ResType =
4073 std::vector<Register> PartialRegs;
4075 unsigned CurrentComponent = 0;
4076 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4080 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4081 TII.get(SPIRV::OpVectorShuffle))
4086 .
addImm(CurrentComponent)
4087 .
addImm(CurrentComponent + 1);
4097 PartialRegs.push_back(SubVecReg);
4100 if (CurrentComponent != ComponentCount) {
4106 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4107 SPIRV::OpVectorExtractDynamic))
4116 PartialRegs.push_back(FinalElemResReg);
4120 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4121 SPIRV::OpCompositeConstruct);
4124bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4125 SPIRVTypeInst ResType,
4129 if (ComponentCount > 2)
4130 return handle64BitOverflow(
4131 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4133 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4135 MachineIRBuilder MIRBuilder(
I);
4139 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4143 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4148 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4155 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4156 TII.get(SPIRV::OpVectorShuffle))
4161 for (
unsigned J = 0; J < ComponentCount; ++J) {
4168 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4171bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4172 SPIRVTypeInst ResType,
4176 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4184bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4185 SPIRVTypeInst ResType,
4186 MachineInstr &
I)
const {
4187 Register OpReg =
I.getOperand(1).getReg();
4196 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4198 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4200 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4202 return SPIRVInstructionSelector::diagnoseUnsupported(
4203 I,
"G_BITREVERSE only support 16,32,64 bits.");
4207 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4218 unsigned AndOp = SPIRV::OpBitwiseAndS;
4219 unsigned OrOp = SPIRV::OpBitwiseOrS;
4220 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4221 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4222 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4224 AndOp = SPIRV::OpBitwiseAndV;
4225 OrOp = SPIRV::OpBitwiseOrV;
4226 ShlOp = SPIRV::OpShiftLeftLogicalV;
4227 ShrOp = SPIRV::OpShiftRightLogicalV;
4233 const unsigned Shift) ->
Register {
4236 (ResType->
getOpcode() != SPIRV::OpTypeVectorIdEXT ||
4243 Register MaskReg = CreateConst(Mask);
4244 Register ShiftReg = CreateConst(Shift);
4251 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4252 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4253 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4254 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4255 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4264 while ((Shift >>= 1) > 0) {
4271 return BuildCOPY(ResVReg, Result,
I);
4274bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4275 SPIRVTypeInst ResType,
4276 MachineInstr &
I)
const {
4277 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4278 "G_FREEZE must define and use a register");
4279 Register OpReg =
I.getOperand(1).getReg();
4283 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4296 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4297 if (
Def->getOpcode() == TargetOpcode::COPY)
4300 switch (
Def->getOpcode()) {
4301 case SPIRV::ASSIGN_TYPE:
4302 if (MachineInstr *AssignToDef =
4304 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4305 Reg =
Def->getOperand(2).getReg();
4308 case SPIRV::OpUndef:
4309 Reg =
Def->getOperand(1).getReg();
4312 unsigned DestOpCode;
4314 DestOpCode = SPIRV::OpConstantNull;
4315 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4316 "static undef/poison lowered to OpConstantNull\n");
4318 DestOpCode = TargetOpcode::COPY;
4320 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4321 "skipped, lowered as a copy of the operand\n");
4323 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4324 .
addDef(
I.getOperand(0).getReg())
4332bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4333 SPIRVTypeInst ResType,
4334 MachineInstr &
I)
const {
4338 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4342 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4347 for (
unsigned i =
I.getNumExplicitDefs();
4348 i <
I.getNumExplicitOperands() && IsConst; ++i)
4353 return diagnoseUnsupported(
4354 I,
"There must be at least two constituent operands in a vector");
4359 for (
unsigned i =
I.getNumExplicitDefs();
4360 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4361 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4366 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4373 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4374 TII.get(IsConst ? SPIRV::OpConstantComposite
4375 : SPIRV::OpCompositeConstruct))
4378 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4379 MIB.
addUse(
I.getOperand(i).getReg());
4384bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4385 SPIRVTypeInst ResType,
4386 MachineInstr &
I)
const {
4390 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4395 unsigned OpIdx =
I.getNumExplicitDefs();
4396 if (!
I.getOperand(OpIdx).isReg())
4400 Register OpReg =
I.getOperand(OpIdx).getReg();
4404 return diagnoseUnsupported(
4405 I,
"There must be at least two constituent operands in a vector");
4408 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4409 TII.get(IsConst ? SPIRV::OpConstantComposite
4410 : SPIRV::OpCompositeConstruct))
4413 for (
unsigned i = 0; i <
N; ++i)
4419bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4420 SPIRVTypeInst ResType,
4421 MachineInstr &
I)
const {
4427 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4429 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4430 TII.get(SPIRV::OpCompositeConstruct))
4433 for (
unsigned OpIdx =
I.getNumExplicitDefs();
4434 OpIdx <
I.getNumExplicitOperands(); ++OpIdx)
4435 MIB.
addUse(
I.getOperand(OpIdx).getReg());
4440bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4441 SPIRVTypeInst ResType,
4442 MachineInstr &
I)
const {
4448 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4450 Opcode = SPIRV::OpDemoteToHelperInvocation;
4452 Opcode = SPIRV::OpKill;
4457 ToErase.eraseFromParent();
4466bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4467 SPIRVTypeInst ResType,
unsigned CmpOpc,
4468 MachineInstr &
I)
const {
4469 Register Cmp0 =
I.getOperand(2).getReg();
4470 Register Cmp1 =
I.getOperand(3).getReg();
4473 "CMP operands should have the same type");
4474 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4484bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4485 SPIRVTypeInst ResType,
4486 MachineInstr &
I)
const {
4487 auto Pred =
I.getOperand(1).getPredicate();
4490 Register CmpOperand =
I.getOperand(2).getReg();
4492 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4497 Register Op1 =
I.getOperand(3).getReg();
4501 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4506 I.getOperand(3).setReg(NewOp1);
4512 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4516SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4517 SPIRVTypeInst ResType)
const {
4519 SPIRVTypeInst SpvI32Ty =
4522 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4529 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4532 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4535 .
addImm(APInt(32, Val).getZExtValue());
4537 GR.
add(ConstInt,
MI);
4544Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4545 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4547 SPIRVTypeInst SpvI32Ty =
4549 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4554 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4555 MachineInstr *
MI =
nullptr;
4559 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4563 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4564 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4570 GR.
add(ConstInt,
MI);
4575bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4576 SPIRVTypeInst ResType,
4577 MachineInstr &
I)
const {
4579 return selectCmp(ResVReg, ResType, CmpOp,
I);
4582bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4583 SPIRVTypeInst ResType,
4584 MachineInstr &
I)
const {
4586 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4596 MachineIRBuilder MIRBuilder(
I);
4603 APFloat ConstVal(3.3219280948873623);
4607 APFloat::rmNearestTiesToEven, &LosesInfo);
4612 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
4614 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4615 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4617 if (!selectExtInst(ResVReg, ResType,
I,
4618 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4628Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4629 MachineInstr &
I)
const {
4637bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4643 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4651 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4654 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4655 Def->getOpcode() == SPIRV::OpConstantI)
4668 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4669 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4671 Intrinsic::spv_const_composite)) {
4672 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4673 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4674 if (!IsZero(
Def->getOperand(i).getReg()))
4683Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4684 MachineInstr &
I)
const {
4693Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4694 MachineInstr &
I)
const {
4704 SPIRVTypeInst ResType,
4705 MachineInstr &
I)
const {
4714bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4715 SPIRVTypeInst ResType,
4716 MachineInstr &
I)
const {
4717 Register SelectFirstArg =
I.getOperand(2).getReg();
4718 Register SelectSecondArg =
I.getOperand(3).getReg();
4732 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4733 }
else if (IsPtrTy) {
4734 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4736 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4739 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4740 "boolean condition");
4742 Opcode = SPIRV::OpSelectSFSCond;
4743 }
else if (IsPtrTy) {
4744 Opcode = SPIRV::OpSelectSPSCond;
4746 Opcode = SPIRV::OpSelectSISCond;
4749 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4752 .
addUse(
I.getOperand(1).getReg())
4761bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4762 SPIRVTypeInst ResType,
4764 MachineInstr &InsertAt,
4765 bool IsSigned)
const {
4767 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4768 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4769 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4771 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4783bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4784 SPIRVTypeInst ResType,
4785 MachineInstr &
I,
bool IsSigned,
4786 unsigned Opcode)
const {
4787 Register SrcReg =
I.getOperand(1).getReg();
4798 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4800 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4803bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4804 SPIRVTypeInst ResType, MachineInstr &
I,
4805 bool IsSigned)
const {
4806 Register SrcReg =
I.getOperand(1).getReg();
4808 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4812 if (ResType == SrcType)
4813 return BuildCOPY(ResVReg, SrcReg,
I);
4815 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4816 return selectUnOp(ResVReg, ResType,
I, Opcode);
4819bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4820 SPIRVTypeInst ResType,
4822 bool IsSigned)
const {
4823 MachineIRBuilder MIRBuilder(
I);
4824 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
4829 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4837 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4840 .
addUse(
I.getOperand(1).getReg())
4841 .
addUse(
I.getOperand(2).getReg())
4846 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4849 .
addUse(
I.getOperand(1).getReg())
4850 .
addUse(
I.getOperand(2).getReg())
4858 unsigned SelectOpcode =
4859 (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4861 ? SPIRV::OpSelectVIVCond
4862 : SPIRV::OpSelectSISCond;
4867 .
addUse(buildOnesVal(
true, ResType,
I))
4868 .
addUse(buildZerosVal(ResType,
I))
4875 .
addUse(buildOnesVal(
false, ResType,
I))
4880bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4883 SPIRVTypeInst IntTy,
4884 SPIRVTypeInst BoolTy)
const {
4888 isVectorType(IntTy) ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4890 Register One = buildOnesVal(
false, IntTy,
I);
4898 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4907bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4908 SPIRVTypeInst ResType,
4909 MachineInstr &
I)
const {
4910 Register IntReg =
I.getOperand(1).getReg();
4913 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4914 if (ArgType == ResType)
4915 return BuildCOPY(ResVReg, IntReg,
I);
4917 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4918 return selectUnOp(ResVReg, ResType,
I, Opcode);
4921bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4922 SPIRVTypeInst ResType,
4923 MachineInstr &
I)
const {
4924 unsigned Opcode =
I.getOpcode();
4925 unsigned TpOpcode = ResType->
getOpcode();
4927 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
4928 assert(Opcode == TargetOpcode::G_CONSTANT &&
4929 I.getOperand(1).getCImm()->isZero());
4930 MachineBasicBlock &DepMBB =
I.getMF()->front();
4933 }
else if (TpOpcode == SPIRV::OpTypeVectorIdEXT) {
4938 "Expected <1 x T> Vector!");
4939 if (Opcode == TargetOpcode::G_FCONSTANT)
4945 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4953 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4956bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4957 SPIRVTypeInst ResType,
4958 MachineInstr &
I)
const {
4959 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4966bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4967 SPIRVTypeInst ResType,
4968 MachineInstr &
I)
const {
4970 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4974 .
addUse(
I.getOperand(3).getReg())
4976 .
addUse(
I.getOperand(2).getReg());
4977 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4983bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4984 SPIRVTypeInst ResType,
4985 MachineInstr &
I)
const {
4986 Type *MaybeResTy =
nullptr;
4991 "Expected aggregate type for extractv instruction");
4993 SPIRV::AccessQualifier::ReadWrite,
false);
4997 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5000 .
addUse(
I.getOperand(2).getReg());
5001 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
5007bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
5008 SPIRVTypeInst ResType,
5009 MachineInstr &
I)
const {
5010 if (
getImm(
I.getOperand(4), MRI))
5011 return selectInsertVal(ResVReg, ResType,
I);
5013 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
5016 .
addUse(
I.getOperand(2).getReg())
5017 .
addUse(
I.getOperand(3).getReg())
5018 .
addUse(
I.getOperand(4).getReg())
5023bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
5024 SPIRVTypeInst ResType,
5025 MachineInstr &
I)
const {
5026 if (
getImm(
I.getOperand(3), MRI))
5027 return selectExtractVal(ResVReg, ResType,
I);
5029 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
5032 .
addUse(
I.getOperand(2).getReg())
5033 .
addUse(
I.getOperand(3).getReg())
5038bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
5039 SPIRVTypeInst ResType,
5040 MachineInstr &
I)
const {
5041 const bool IsGEPInBounds =
I.getOperand(2).getImm();
5044 const bool UseUntypedPointers =
5045 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
5050 if (UseUntypedPointers) {
5052 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
5053 : SPIRV::OpUntypedAccessChainKHR;
5055 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
5056 : SPIRV::OpUntypedPtrAccessChainKHR;
5065 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
5067 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
5068 : SPIRV::OpPtrAccessChain;
5073 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5078 if (UseUntypedPointers) {
5093 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5094 Def->getOperand(1).isReg())
5096 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5097 if (
const auto *GVar =
5100 SPIRV::AccessQualifier::ReadWrite,
5104 return diagnoseUnsupported(
5105 I,
"could not deduce the base type of an untyped access chain");
5110 Res.addUse(BaseReg);
5112 const bool IsAccessChainOpcode =
5113 (Opcode == SPIRV::OpAccessChain ||
5114 Opcode == SPIRV::OpInBoundsAccessChain ||
5115 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5116 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5118 assert((!IsAccessChainOpcode || (
getImm(
I.getOperand(4), MRI) &&
5119 foldImm(
I.getOperand(4), MRI) == 0)) &&
5120 "Cannot translate GEP to OpAccessChain.");
5123 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5124 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5125 Res.addUse(
I.getOperand(i).getReg());
5126 Res.constrainAllUses(
TII,
TRI, RBI);
5135 if (Extract.
getOpcode() == SPIRV::OpCompositeExtract) {
5141 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5147 TII.get(SPIRV::OpCompositeInsert))
5161bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5163 unsigned Lim =
I.getNumExplicitOperands();
5164 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5165 Register OpReg =
I.getOperand(i).getReg();
5166 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
5168 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
5169 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5170 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5183 SPIRVTypeInst WrapType = OpType;
5184 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5186 SPIRV::StorageClass::CodeSectionINTEL) {
5188 SPIRV::StorageClass::Function,
I);
5195 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5196 TII.get(SPIRV::OpSpecConstantOp))
5199 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5201 GR.
add(OpDefine, MIB);
5207bool SPIRVInstructionSelector::selectDerivativeInst(
5208 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5209 const unsigned DPdOpCode)
const {
5212 if (!errorIfInstrOutsideShader(
I))
5218 Register SrcReg =
I.getOperand(2).getReg();
5223 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5226 .
addUse(
I.getOperand(2).getReg());
5228 MachineIRBuilder MIRBuilder(
I);
5231 if (componentCount != 1)
5239 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5244 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5249 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5257bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5258 SPIRVTypeInst ResType,
5259 MachineInstr &
I)
const {
5263 case Intrinsic::spv_load:
5264 return selectLoad(ResVReg, ResType,
I);
5265 case Intrinsic::spv_atomic_load:
5266 return selectAtomicLoad(ResVReg, ResType,
I);
5267 case Intrinsic::spv_store:
5268 return selectStore(
I);
5269 case Intrinsic::spv_atomic_store:
5270 return selectAtomicStore(
I);
5271 case Intrinsic::spv_extractv:
5272 return selectExtractVal(ResVReg, ResType,
I);
5273 case Intrinsic::spv_insertv:
5274 return selectInsertVal(ResVReg, ResType,
I);
5275 case Intrinsic::spv_extractelt:
5276 return selectExtractElt(ResVReg, ResType,
I);
5277 case Intrinsic::spv_insertelt:
5278 return selectInsertElt(ResVReg, ResType,
I);
5279 case Intrinsic::spv_gep:
5280 return selectGEP(ResVReg, ResType,
I);
5281 case Intrinsic::spv_bitcast: {
5282 Register OpReg =
I.getOperand(2).getReg();
5283 SPIRVTypeInst OpType =
5287 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5289 case Intrinsic::spv_unref_global:
5290 case Intrinsic::spv_init_global: {
5291 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
5296 Register GVarVReg =
MI->getOperand(0).getReg();
5297 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5302 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5304 MI->eraseFromParent();
5308 case Intrinsic::spv_undef: {
5309 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5315 case Intrinsic::spv_poison:
5316 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5321 case Intrinsic::spv_freeze:
5322 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5325 .
addUse(
I.getOperand(2).getReg())
5328 case Intrinsic::spv_named_boolean_spec_constant: {
5329 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5330 : SPIRV::OpSpecConstantFalse;
5332 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5333 .
addDef(
I.getOperand(0).getReg())
5336 unsigned SpecId =
I.getOperand(2).getImm();
5338 SPIRV::Decoration::SpecId, {SpecId});
5342 case Intrinsic::spv_const_composite: {
5344 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5350 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5352 std::function<bool(
Register)> HasSpecConstOperand =
5362 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5363 J < Def->getNumExplicitOperands(); ++J) {
5364 if (
Def->getOperand(J).isReg() &&
5365 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5371 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5372 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5373 : SPIRV::OpConstantComposite;
5374 unsigned ContinuedOpc = HasSpecConst
5375 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5376 : SPIRV::OpConstantCompositeContinuedINTEL;
5377 MachineIRBuilder MIR(
I);
5379 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5381 for (
auto *Instr : Instructions) {
5382 Instr->setDebugLoc(
I.getDebugLoc());
5387 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5394 case Intrinsic::spv_assign_name: {
5395 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5396 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5397 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5398 i <
I.getNumExplicitOperands(); ++i) {
5399 MIB.
addImm(
I.getOperand(i).getImm());
5404 case Intrinsic::spv_switch: {
5405 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5406 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5407 if (
I.getOperand(i).isReg())
5408 MIB.
addReg(
I.getOperand(i).getReg());
5409 else if (
I.getOperand(i).isCImm())
5410 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5411 else if (
I.getOperand(i).isMBB())
5412 MIB.
addMBB(
I.getOperand(i).getMBB());
5419 case Intrinsic::spv_loop_merge: {
5420 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5421 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5422 if (
I.getOperand(i).isMBB())
5423 MIB.
addMBB(
I.getOperand(i).getMBB());
5430 case Intrinsic::spv_loop_control_intel: {
5432 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5433 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5438 case Intrinsic::spv_selection_merge: {
5440 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5441 assert(
I.getOperand(1).isMBB() &&
5442 "operand 1 to spv_selection_merge must be a basic block");
5443 MIB.
addMBB(
I.getOperand(1).getMBB());
5444 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5448 case Intrinsic::spv_cmpxchg:
5449 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5450 case Intrinsic::spv_unreachable:
5451 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5454 case Intrinsic::spv_abort:
5455 return selectAbort(
I);
5456 case Intrinsic::spv_alloca:
5457 return selectFrameIndex(ResVReg, ResType,
I);
5458 case Intrinsic::spv_alloca_array:
5459 return selectAllocaArray(ResVReg, ResType,
I);
5460 case Intrinsic::spv_assume:
5462 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5463 .
addUse(
I.getOperand(1).getReg())
5468 case Intrinsic::spv_expect:
5470 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5473 .
addUse(
I.getOperand(2).getReg())
5474 .
addUse(
I.getOperand(3).getReg())
5479 case Intrinsic::arithmetic_fence:
5480 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5481 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5484 .
addUse(
I.getOperand(2).getReg())
5488 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5490 case Intrinsic::spv_thread_id:
5496 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5498 case Intrinsic::spv_thread_id_in_group:
5504 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5506 case Intrinsic::spv_group_id:
5512 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5514 case Intrinsic::spv_flattened_thread_id_in_group:
5521 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5523 case Intrinsic::spv_workgroup_size:
5524 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5526 case Intrinsic::spv_global_size:
5527 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5529 case Intrinsic::spv_global_offset:
5530 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5532 case Intrinsic::spv_num_workgroups:
5533 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5535 case Intrinsic::spv_subgroup_size:
5536 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5538 case Intrinsic::spv_num_subgroups:
5539 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5541 case Intrinsic::spv_subgroup_id:
5542 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5543 case Intrinsic::spv_subgroup_local_invocation_id:
5544 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5545 ResVReg, ResType,
I);
5546 case Intrinsic::spv_subgroup_max_size:
5547 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5549 case Intrinsic::spv_fdot:
5550 return selectFloatDot(ResVReg, ResType,
I);
5551 case Intrinsic::spv_udot:
5552 case Intrinsic::spv_sdot:
5553 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5555 return selectIntegerDot(ResVReg, ResType,
I,
5556 IID == Intrinsic::spv_sdot);
5557 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5558 case Intrinsic::spv_dot4add_i8packed:
5559 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5561 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5562 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5563 case Intrinsic::spv_dot4add_u8packed:
5564 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5566 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5567 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5568 case Intrinsic::spv_all:
5569 return selectAll(ResVReg, ResType,
I);
5570 case Intrinsic::spv_any:
5571 return selectAny(ResVReg, ResType,
I);
5572 case Intrinsic::spv_distance:
5573 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5574 case Intrinsic::spv_lerp:
5575 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5576 case Intrinsic::spv_length:
5577 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5578 case Intrinsic::spv_degrees:
5579 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5580 case Intrinsic::spv_faceforward:
5581 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5582 case Intrinsic::spv_frac:
5583 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5584 case Intrinsic::spv_isinf:
5585 return selectOpIsInf(ResVReg, ResType,
I);
5586 case Intrinsic::spv_isnan:
5587 return selectOpIsNan(ResVReg, ResType,
I);
5588 case Intrinsic::spv_isfinite:
5589 return selectOpIsFinite(ResVReg, ResType,
I);
5590 case Intrinsic::spv_isnormal:
5591 return selectOpIsNormal(ResVReg, ResType,
I);
5592 case Intrinsic::spv_normalize:
5593 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5594 case Intrinsic::spv_refract:
5595 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5596 case Intrinsic::spv_reflect:
5597 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5598 case Intrinsic::spv_rsqrt:
5599 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5600 case Intrinsic::spv_sign:
5601 return selectSign(ResVReg, ResType,
I);
5602 case Intrinsic::spv_smoothstep:
5603 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5604 case Intrinsic::spv_firstbituhigh:
5605 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5606 case Intrinsic::spv_firstbitshigh:
5607 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5608 case Intrinsic::spv_firstbitlow:
5609 return selectFirstBitLow(ResVReg, ResType,
I);
5610 case Intrinsic::spv_all_memory_barrier:
5611 return selectBarrierInst(
I, SPIRV::Scope::Device,
5612 SPIRV::MemorySemantics::UniformMemory |
5613 SPIRV::MemorySemantics::ImageMemory |
5614 SPIRV::MemorySemantics::WorkgroupMemory,
5616 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5617 return selectBarrierInst(
I, SPIRV::Scope::Device,
5618 SPIRV::MemorySemantics::UniformMemory |
5619 SPIRV::MemorySemantics::ImageMemory |
5620 SPIRV::MemorySemantics::WorkgroupMemory,
5622 case Intrinsic::spv_device_memory_barrier:
5623 return selectBarrierInst(
I, SPIRV::Scope::Device,
5624 SPIRV::MemorySemantics::UniformMemory |
5625 SPIRV::MemorySemantics::ImageMemory,
5627 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5628 return selectBarrierInst(
I, SPIRV::Scope::Device,
5629 SPIRV::MemorySemantics::UniformMemory |
5630 SPIRV::MemorySemantics::ImageMemory,
5632 case Intrinsic::spv_group_memory_barrier:
5633 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5634 SPIRV::MemorySemantics::WorkgroupMemory,
5636 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5637 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5638 SPIRV::MemorySemantics::WorkgroupMemory,
5640 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5641 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5642 SPIRV::StorageClass::StorageClass ResSC =
5645 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5646 "from the Generic storage class");
5647 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5655 case Intrinsic::spv_lifetime_start:
5656 case Intrinsic::spv_lifetime_end: {
5657 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5658 : SPIRV::OpLifetimeStop;
5659 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5660 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5669 case Intrinsic::spv_saturate:
5670 return selectSaturate(ResVReg, ResType,
I);
5671 case Intrinsic::spv_nclamp:
5672 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5673 case Intrinsic::spv_uclamp:
5674 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5675 case Intrinsic::spv_sclamp:
5676 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5677 case Intrinsic::spv_subgroup_prefix_bit_count:
5678 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5679 case Intrinsic::spv_wave_active_countbits:
5680 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5681 case Intrinsic::spv_wave_all_equal:
5682 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5683 case Intrinsic::spv_wave_all:
5684 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5685 case Intrinsic::spv_wave_any:
5686 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5687 case Intrinsic::spv_subgroup_ballot:
5688 return selectWaveOpInst(ResVReg, ResType,
I,
5689 SPIRV::OpGroupNonUniformBallot);
5690 case Intrinsic::spv_wave_is_first_lane:
5691 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5692 case Intrinsic::spv_wave_reduce_or:
5693 return selectWaveReduceOp(ResVReg, ResType,
I,
5694 SPIRV::OpGroupNonUniformBitwiseOr);
5695 case Intrinsic::spv_wave_reduce_xor:
5696 return selectWaveReduceOp(ResVReg, ResType,
I,
5697 SPIRV::OpGroupNonUniformBitwiseXor);
5698 case Intrinsic::spv_wave_reduce_and:
5699 return selectWaveReduceOp(ResVReg, ResType,
I,
5700 SPIRV::OpGroupNonUniformBitwiseAnd);
5701 case Intrinsic::spv_wave_reduce_umax:
5702 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5703 case Intrinsic::spv_wave_reduce_max:
5704 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5705 case Intrinsic::spv_wave_reduce_umin:
5706 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5707 case Intrinsic::spv_wave_reduce_min:
5708 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5709 case Intrinsic::spv_wave_reduce_sum:
5710 return selectWaveReduceSum(ResVReg, ResType,
I);
5711 case Intrinsic::spv_wave_product:
5712 return selectWaveReduceProduct(ResVReg, ResType,
I);
5713 case Intrinsic::spv_wave_readlane:
5714 return selectWaveOpInst(ResVReg, ResType,
I,
5715 SPIRV::OpGroupNonUniformShuffle);
5716 case Intrinsic::spv_wave_prefix_sum:
5717 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5718 case Intrinsic::spv_wave_prefix_product:
5719 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5720 case Intrinsic::spv_quad_read_across_x: {
5721 return selectQuadSwap(ResVReg, ResType,
I, 0);
5723 case Intrinsic::spv_quad_read_across_y: {
5724 return selectQuadSwap(ResVReg, ResType,
I, 1);
5726 case Intrinsic::spv_quad_read_across_diagonal: {
5727 return selectQuadSwap(ResVReg, ResType,
I, 2);
5729 case Intrinsic::spv_radians:
5730 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5734 case Intrinsic::instrprof_increment:
5735 case Intrinsic::instrprof_increment_step:
5736 case Intrinsic::instrprof_value_profile:
5739 case Intrinsic::spv_value_md:
5741 case Intrinsic::spv_resource_handlefrombinding: {
5742 return selectHandleFromBinding(ResVReg, ResType,
I);
5744 case Intrinsic::spv_resource_counterhandlefrombinding:
5745 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5746 case Intrinsic::spv_resource_updatecounter:
5747 return selectUpdateCounter(ResVReg, ResType,
I);
5748 case Intrinsic::spv_resource_store_typedbuffer: {
5749 return selectImageWriteIntrinsic(
I);
5751 case Intrinsic::spv_resource_load_typedbuffer: {
5752 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5754 case Intrinsic::spv_resource_load_level: {
5755 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5757 case Intrinsic::spv_resource_getdimensions_x:
5758 case Intrinsic::spv_resource_getdimensions_xy:
5759 case Intrinsic::spv_resource_getdimensions_xyz: {
5760 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5762 case Intrinsic::spv_resource_getdimensions_levels_x:
5763 case Intrinsic::spv_resource_getdimensions_levels_xy:
5764 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5765 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5767 case Intrinsic::spv_resource_getdimensions_ms_xy:
5768 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5769 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5771 case Intrinsic::spv_resource_calculate_lod:
5772 case Intrinsic::spv_resource_calculate_lod_unclamped:
5773 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5774 case Intrinsic::spv_resource_sample:
5775 case Intrinsic::spv_resource_sample_clamp:
5776 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5777 case Intrinsic::spv_resource_samplebias:
5778 case Intrinsic::spv_resource_samplebias_clamp:
5779 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5780 case Intrinsic::spv_resource_samplegrad:
5781 case Intrinsic::spv_resource_samplegrad_clamp:
5782 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5783 case Intrinsic::spv_resource_samplelevel:
5784 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5785 case Intrinsic::spv_resource_samplecmp:
5786 case Intrinsic::spv_resource_samplecmp_clamp:
5787 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5788 case Intrinsic::spv_resource_samplecmplevelzero:
5789 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5790 case Intrinsic::spv_resource_gather:
5791 case Intrinsic::spv_resource_gather_cmp:
5792 return selectGatherIntrinsic(ResVReg, ResType,
I);
5793 case Intrinsic::spv_resource_getbasepointer:
5794 case Intrinsic::spv_resource_getpointer: {
5795 return selectResourceGetPointer(ResVReg, ResType,
I);
5797 case Intrinsic::spv_pushconstant_getpointer: {
5798 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5800 case Intrinsic::spv_discard: {
5801 return selectDiscard(ResVReg, ResType,
I);
5803 case Intrinsic::spv_resource_nonuniformindex: {
5804 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5806 case Intrinsic::spv_unpackhalf2x16: {
5807 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5809 case Intrinsic::spv_packhalf2x16: {
5810 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5812 case Intrinsic::spv_ddx:
5813 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5814 case Intrinsic::spv_ddy:
5815 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5816 case Intrinsic::spv_ddx_coarse:
5817 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5818 case Intrinsic::spv_ddy_coarse:
5819 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5820 case Intrinsic::spv_ddx_fine:
5821 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5822 case Intrinsic::spv_ddy_fine:
5823 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5824 case Intrinsic::spv_fwidth:
5825 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5826 case Intrinsic::spv_masked_gather:
5827 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5828 return selectMaskedGather(ResVReg, ResType,
I);
5829 return diagnoseUnsupported(
5830 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5831 case Intrinsic::spv_masked_scatter:
5832 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5833 return selectMaskedScatter(
I);
5834 return diagnoseUnsupported(
5835 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5836 case Intrinsic::returnaddress:
5837 case Intrinsic::frameaddress: {
5839 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5846 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5851bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5852 SPIRVTypeInst ResType,
5853 MachineInstr &
I)
const {
5856 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5863bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5864 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5866 assert(Intr.getIntrinsicID() ==
5867 Intrinsic::spv_resource_counterhandlefrombinding);
5870 Register MainHandleReg = Intr.getOperand(2).getReg();
5872 assert(MainHandleDef->getIntrinsicID() ==
5873 Intrinsic::spv_resource_handlefrombinding);
5877 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5878 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5879 std::string CounterName =
5884 MachineIRBuilder MIRBuilder(
I);
5886 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5888 ArraySize, IndexReg, CounterName, MIRBuilder);
5890 return BuildCOPY(ResVReg, CounterVarReg,
I);
5893bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5894 SPIRVTypeInst ResType,
5895 MachineInstr &
I)
const {
5897 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5899 Register CounterHandleReg = Intr.getOperand(2).getReg();
5900 Register IncrReg = Intr.getOperand(3).getReg();
5907 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5908 assert(CounterVarPointeeType &&
5909 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5910 "Counter variable must be a struct");
5912 SPIRV::StorageClass::StorageBuffer &&
5913 "Counter variable must be in the storage buffer storage class");
5915 "Counter variable must have exactly 1 member in the struct");
5916 const SPIRVTypeInst MemberType =
5919 "Counter variable struct must have a single i32 member");
5923 MachineIRBuilder MIRBuilder(
I);
5925 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5928 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5934 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5937 .
addUse(CounterHandleReg)
5944 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5947 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5950 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5959 return BuildCOPY(ResVReg, AtomicRes,
I);
5967 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5975bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5976 SPIRVTypeInst ResType,
5977 MachineInstr &
I)
const {
5985 Register ImageReg =
I.getOperand(2).getReg();
5993 Register IdxReg =
I.getOperand(3).getReg();
5995 MachineInstr &Pos =
I;
5997 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
6001bool SPIRVInstructionSelector::generateSampleImage(
6004 DebugLoc Loc, MachineInstr &Pos)
const {
6015 if (!loadHandleBeforePosition(NewSamplerReg,
6021 MachineIRBuilder MIRBuilder(Pos);
6034 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
6035 ImOps.Lod.has_value();
6036 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
6037 : SPIRV::OpImageSampleImplicitLod;
6039 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
6040 : SPIRV::OpImageSampleDrefImplicitLod;
6049 MIB.
addUse(*ImOps.Compare);
6051 uint32_t ImageOperands = 0;
6053 ImageOperands |= SPIRV::ImageOperand::Bias;
6055 ImageOperands |= SPIRV::ImageOperand::Lod;
6056 if (ImOps.GradX && ImOps.GradY)
6057 ImageOperands |= SPIRV::ImageOperand::Grad;
6058 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
6060 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6063 "Non-constant offsets are not supported in sample instructions.");
6068 ImageOperands |= SPIRV::ImageOperand::MinLod;
6070 if (ImageOperands != 0) {
6071 MIB.
addImm(ImageOperands);
6072 if (ImageOperands & SPIRV::ImageOperand::Bias)
6074 if (ImageOperands & SPIRV::ImageOperand::Lod)
6076 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6077 MIB.
addUse(*ImOps.GradX);
6078 MIB.
addUse(*ImOps.GradY);
6081 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6082 MIB.
addUse(*ImOps.Offset);
6083 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6084 MIB.
addUse(*ImOps.MinLod);
6091bool SPIRVInstructionSelector::selectImageQuerySize(
6093 std::optional<Register> LodReg)
const {
6095 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6098 "ImageReg is not an image type.");
6100 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6102 unsigned NumComponents = 0;
6104 case SPIRV::Dim::DIM_1D:
6105 case SPIRV::Dim::DIM_Buffer:
6106 NumComponents =
IsArray ? 2 : 1;
6108 case SPIRV::Dim::DIM_2D:
6109 case SPIRV::Dim::DIM_Cube:
6110 case SPIRV::Dim::DIM_Rect:
6111 NumComponents =
IsArray ? 3 : 2;
6113 case SPIRV::Dim::DIM_3D:
6117 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6122 SPIRVTypeInst ResType =
6127 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6137bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6138 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6139 Register ImageReg =
I.getOperand(2).getReg();
6146 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6149bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6150 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6151 Register ImageReg =
I.getOperand(2).getReg();
6160 Register LodReg =
I.getOperand(3).getReg();
6163 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6165 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6172 TII.get(SPIRV::OpImageQueryLevels))
6179 TII.get(SPIRV::OpCompositeConstruct))
6189bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6190 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6191 Register ImageReg =
I.getOperand(2).getReg();
6202 "OpImageQuerySamples requires a multisampled image");
6204 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6212 TII.get(SPIRV::OpImageQuerySamples))
6219 TII.get(SPIRV::OpCompositeConstruct))
6229bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6230 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6231 Register ImageReg =
I.getOperand(2).getReg();
6232 Register SamplerReg =
I.getOperand(3).getReg();
6233 Register CoordinateReg =
I.getOperand(4).getReg();
6249 if (!loadHandleBeforePosition(
6254 MachineIRBuilder MIRBuilder(
I);
6260 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6270 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6277 unsigned ExtractedIndex =
6279 Intrinsic::spv_resource_calculate_lod_unclamped
6283 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6284 TII.get(SPIRV::OpCompositeExtract))
6294bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6295 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6296 Register ImageReg =
I.getOperand(2).getReg();
6297 Register SamplerReg =
I.getOperand(3).getReg();
6298 Register CoordinateReg =
I.getOperand(4).getReg();
6299 ImageOperands ImOps;
6300 if (
I.getNumOperands() > 5)
6301 ImOps.Offset =
I.getOperand(5).getReg();
6302 if (
I.getNumOperands() > 6)
6303 ImOps.MinLod =
I.getOperand(6).getReg();
6304 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6305 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6308bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6309 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6310 Register ImageReg =
I.getOperand(2).getReg();
6311 Register SamplerReg =
I.getOperand(3).getReg();
6312 Register CoordinateReg =
I.getOperand(4).getReg();
6313 ImageOperands ImOps;
6314 ImOps.Bias =
I.getOperand(5).getReg();
6315 if (
I.getNumOperands() > 6)
6316 ImOps.Offset =
I.getOperand(6).getReg();
6317 if (
I.getNumOperands() > 7)
6318 ImOps.MinLod =
I.getOperand(7).getReg();
6319 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6320 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6323bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6324 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6325 Register ImageReg =
I.getOperand(2).getReg();
6326 Register SamplerReg =
I.getOperand(3).getReg();
6327 Register CoordinateReg =
I.getOperand(4).getReg();
6328 ImageOperands ImOps;
6329 ImOps.GradX =
I.getOperand(5).getReg();
6330 ImOps.GradY =
I.getOperand(6).getReg();
6331 if (
I.getNumOperands() > 7)
6332 ImOps.Offset =
I.getOperand(7).getReg();
6333 if (
I.getNumOperands() > 8)
6334 ImOps.MinLod =
I.getOperand(8).getReg();
6335 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6336 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6339bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6340 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6341 Register ImageReg =
I.getOperand(2).getReg();
6342 Register SamplerReg =
I.getOperand(3).getReg();
6343 Register CoordinateReg =
I.getOperand(4).getReg();
6344 ImageOperands ImOps;
6345 ImOps.Lod =
I.getOperand(5).getReg();
6346 if (
I.getNumOperands() > 6)
6347 ImOps.Offset =
I.getOperand(6).getReg();
6348 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6349 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6352bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6353 SPIRVTypeInst ResType,
6354 MachineInstr &
I)
const {
6355 Register ImageReg =
I.getOperand(2).getReg();
6356 Register SamplerReg =
I.getOperand(3).getReg();
6357 Register CoordinateReg =
I.getOperand(4).getReg();
6358 ImageOperands ImOps;
6359 ImOps.Compare =
I.getOperand(5).getReg();
6360 if (
I.getNumOperands() > 6)
6361 ImOps.Offset =
I.getOperand(6).getReg();
6362 if (
I.getNumOperands() > 7)
6363 ImOps.MinLod =
I.getOperand(7).getReg();
6364 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6365 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6368bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6369 SPIRVTypeInst ResType,
6370 MachineInstr &
I)
const {
6371 Register ImageReg =
I.getOperand(2).getReg();
6372 Register CoordinateReg =
I.getOperand(3).getReg();
6373 Register LodReg =
I.getOperand(4).getReg();
6375 ImageOperands ImOps;
6377 if (
I.getNumOperands() > 5)
6378 ImOps.Offset =
I.getOperand(5).getReg();
6390 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6391 I.getDebugLoc(),
I, &ImOps);
6394bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6395 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6396 Register ImageReg =
I.getOperand(2).getReg();
6397 Register SamplerReg =
I.getOperand(3).getReg();
6398 Register CoordinateReg =
I.getOperand(4).getReg();
6399 ImageOperands ImOps;
6400 ImOps.Compare =
I.getOperand(5).getReg();
6401 if (
I.getNumOperands() > 6)
6402 ImOps.Offset =
I.getOperand(6).getReg();
6405 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6406 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6409bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6410 SPIRVTypeInst ResType,
6411 MachineInstr &
I)
const {
6412 Register ImageReg =
I.getOperand(2).getReg();
6413 Register SamplerReg =
I.getOperand(3).getReg();
6414 Register CoordinateReg =
I.getOperand(4).getReg();
6417 "ImageReg is not an image type.");
6422 ComponentOrCompareReg =
I.getOperand(5).getReg();
6423 OffsetReg =
I.getOperand(6).getReg();
6426 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6430 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6431 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6432 Dim != SPIRV::Dim::DIM_Rect) {
6434 "Gather operations are only supported for 2D, Cube, and Rect images.");
6441 if (!loadHandleBeforePosition(
6446 MachineIRBuilder MIRBuilder(
I);
6447 SPIRVTypeInst SampledImageType =
6452 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6460 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6462 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6464 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6469 .
addUse(ComponentOrCompareReg);
6471 uint32_t ImageOperands = 0;
6472 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6473 if (Dim == SPIRV::Dim::DIM_Cube) {
6475 "Gather operations with offset are not supported for Cube images.");
6479 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6481 ImageOperands |= SPIRV::ImageOperand::Offset;
6485 if (ImageOperands != 0) {
6486 MIB.
addImm(ImageOperands);
6488 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6496bool SPIRVInstructionSelector::generateImageReadOrFetch(
6499 const ImageOperands *ImOps)
const {
6502 "ImageReg is not an image type.");
6504 bool IsSignedInteger =
6509 bool IsFetch = (SampledOp.getImm() == 1);
6511 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6512 uint32_t ImageOperandsMask = 0;
6513 if (IsSignedInteger)
6514 ImageOperandsMask |= 0x1000;
6516 if (IsFetch && ImOps) {
6518 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6519 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6521 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6523 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6527 if (ImageOperandsMask != 0) {
6528 MIB.
addImm(ImageOperandsMask);
6529 if (IsFetch && ImOps) {
6532 if (ImOps->Offset &&
6533 (ImageOperandsMask &
6534 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6535 MIB.
addUse(*ImOps->Offset);
6544 SPIRVTypeInst SampledType =
6547 SPIRVTypeInst ReadType =
6548 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6549 bool ReadTypeMatchesResult = ReadType == ResType;
6551 Register ReadReg = ReadTypeMatchesResult
6557 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6563 BMI.constrainAllUses(
TII,
TRI, RBI);
6565 if (ReadTypeMatchesResult)
6578 if (ResultSize == 1) {
6587 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6590bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6591 SPIRVTypeInst ResType,
6592 MachineInstr &
I)
const {
6593 Register ResourcePtr =
I.getOperand(2).getReg();
6595 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6604 MachineIRBuilder MIRBuilder(
I);
6609 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6615 if (
I.getNumExplicitOperands() > 3) {
6616 Register IndexReg =
I.getOperand(3).getReg();
6623bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6624 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6629bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6630 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6631 Register ObjReg =
I.getOperand(2).getReg();
6632 if (!BuildCOPY(ResVReg, ObjReg,
I))
6642 decorateUsesAsNonUniform(ResVReg);
6646void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6649 {NonUniformReg,
nullptr}};
6650 llvm::SmallSet<Register, 8> Visited;
6651 while (WorkList.
size() > 0) {
6654 if (!Visited.
insert(CurrentReg).second)
6657 bool IsDecorated =
false;
6659 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6660 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6666 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6668 if (ResultReg == CurrentReg)
6676 MachineInstr &InsertPt =
6679 SPIRV::Decoration::NonUniformEXT, {});
6684bool SPIRVInstructionSelector::extractSubvector(
6686 MachineInstr &InsertionPoint)
const {
6688 [[maybe_unused]]
uint64_t InputSize =
6691 [[maybe_unused]]
bool IsLongVectorEXT =
6693 assert((InputSize > 1 || IsLongVectorEXT) &&
"The input must be a vector.");
6694 assert((ResultSize > 1 || IsLongVectorEXT) &&
"The result must be a vector.");
6695 assert(ResultSize < InputSize &&
6696 "Cannot extract more element than there are in the input.");
6703 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6712 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6714 TII.get(SPIRV::OpCompositeConstruct))
6718 for (
Register ComponentReg : ComponentRegisters)
6719 MIB.
addUse(ComponentReg);
6724bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6725 MachineInstr &
I)
const {
6732 Register ImageReg =
I.getOperand(1).getReg();
6740 Register CoordinateReg =
I.getOperand(2).getReg();
6741 Register DataReg =
I.getOperand(3).getReg();
6744 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6752Register SPIRVInstructionSelector::buildPointerToResource(
6753 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6754 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6755 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6757 if (ArraySize == 1) {
6758 SPIRVTypeInst PtrType =
6761 "SpirvResType did not have an explicit layout.");
6766 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6767 SPIRVTypeInst VarPointerType =
6770 VarPointerType, Set,
Binding, Name, MIRBuilder);
6772 SPIRVTypeInst ResPointerType =
6785bool SPIRVInstructionSelector::selectFirstBitSet16(
6786 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6787 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6789 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6793 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6796bool SPIRVInstructionSelector::selectFirstBitSet32(
6798 unsigned BitSetOpcode)
const {
6799 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6802 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6809bool SPIRVInstructionSelector::selectFirstBitSet64(
6811 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6825 if (ComponentCount > 2) {
6826 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6828 unsigned Opcode) ->
bool {
6829 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6833 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6837 MachineIRBuilder MIRBuilder(
I);
6839 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6843 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6849 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6859 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6860 SPIRV::OpVectorExtractDynamic))
6862 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6863 SPIRV::OpVectorExtractDynamic))
6867 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6868 TII.get(SPIRV::OpVectorShuffle))
6876 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6882 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6883 TII.get(SPIRV::OpVectorShuffle))
6891 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6911 SelectOp = SPIRV::OpSelectSISCond;
6912 AddOp = SPIRV::OpIAddS;
6920 SelectOp = SPIRV::OpSelectVIVCond;
6921 AddOp = SPIRV::OpIAddV;
6927 Register RegSecondaryOffset = Reg0;
6931 if (SwapPrimarySide) {
6932 PrimaryReg = LowReg;
6933 SecondaryReg = HighReg;
6934 RegPrimaryOffset = Reg0;
6935 RegSecondaryOffset = Reg32;
6940 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6941 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6946 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6947 SPIRV::OpINotEqual))
6954 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6955 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6960 if (SwapPrimarySide) {
6962 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6963 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6974 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6975 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6980 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6981 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6984 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6988bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6989 SPIRVTypeInst ResType,
6991 bool IsSigned)
const {
6993 Register OpReg =
I.getOperand(2).getReg();
6996 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6997 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
7001 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7003 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7005 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7008 return diagnoseUnsupported(
7010 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
7014bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
7015 SPIRVTypeInst ResType,
7016 MachineInstr &
I)
const {
7018 Register OpReg =
I.getOperand(2).getReg();
7023 unsigned ExtendOpcode = SPIRV::OpUConvert;
7024 unsigned BitSetOpcode = GL::FindILsb;
7028 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7030 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7032 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7035 return diagnoseUnsupported(
I,
7036 "spv_firstbitlow only supports 16,32,64 bits.");
7040bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
7041 SPIRVTypeInst ResType,
7042 MachineInstr &
I)
const {
7046 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
7049 .
addUse(
I.getOperand(2).getReg())
7052 unsigned Alignment =
I.getOperand(3).getImm();
7066 while (!Worklist.
empty()) {
7068 switch (
T->getOpcode()) {
7069 case SPIRV::OpTypeInt:
7070 case SPIRV::OpTypeFloat:
7071 case SPIRV::OpTypePointer:
7073 case SPIRV::OpTypeVector:
7074 case SPIRV::OpTypeVectorIdEXT:
7075 case SPIRV::OpTypeMatrix:
7076 case SPIRV::OpTypeArray: {
7077 Register OperandReg =
T->getOperand(1).getReg();
7081 case SPIRV::OpTypeStruct:
7082 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7083 Register OperandReg =
T->getOperand(Idx).getReg();
7095bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7096 assert(
I.getNumExplicitOperands() == 2);
7098 Register MsgReg =
I.getOperand(1).getReg();
7100 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7103 return diagnoseUnsupported(
7105 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7106 "scalar, pointer, vector, matrix, or aggregate of such types)");
7109 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7116bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7125 uint32_t MsgVal = ~0
u;
7126 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7127 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7130 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7133 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7140bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7141 SPIRVTypeInst ResType,
7142 MachineInstr &
I)
const {
7149 bool UseUntypedPointers =
7150 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7152 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7154 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7157 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7161 if (UseUntypedPointers) {
7165 return diagnoseUnsupported(
7166 I,
"could not deduce the data type of an untyped variable");
7172 unsigned Alignment =
I.getOperand(2).getImm();
7179bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7184 const MachineInstr *PrevI =
I.getPrevNode();
7186 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7190 .
addMBB(
I.getOperand(0).getMBB())
7195 .
addMBB(
I.getOperand(0).getMBB())
7200bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7211 const MachineInstr *NextI =
I.getNextNode();
7213 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7219 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7221 .
addUse(
I.getOperand(0).getReg())
7222 .
addMBB(
I.getOperand(1).getMBB())
7228bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7229 MachineInstr &
I)
const {
7231 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7233 const unsigned NumOps =
I.getNumOperands();
7234 for (
unsigned i = 1; i <
NumOps; i += 2) {
7235 MIB.
addUse(
I.getOperand(i + 0).getReg());
7236 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7242bool SPIRVInstructionSelector::selectGlobalValue(
7243 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7245 MachineIRBuilder MIRBuilder(
I);
7246 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7249 std::string GlobalIdent;
7251 unsigned &
ID = UnnamedGlobalIDs[GV];
7253 ID = UnnamedGlobalIDs.
size();
7254 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7280 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7287 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7292 MachineInstrBuilder MIB1 =
7293 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7296 MachineInstrBuilder MIB2 =
7298 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7302 GR.
add(ConstVal, MIB2);
7310 MachineInstrBuilder MIB3 =
7311 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7314 GR.
add(ConstVal, MIB3);
7320 assert(NewReg != ResVReg);
7321 return BuildCOPY(ResVReg, NewReg,
I);
7331 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7334 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7340 SPIRVTypeInst ResType =
7344 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7349 if (
GlobalVar->isExternallyInitialized() &&
7350 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7351 constexpr unsigned ReadWriteINTEL = 3u;
7354 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7360bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7361 SPIRVTypeInst ResType,
7362 MachineInstr &
I)
const {
7364 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7372 MachineIRBuilder MIRBuilder(
I);
7377 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7380 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7382 .
add(
I.getOperand(1))
7396 APFloat::rmNearestTiesToEven, &LosesInfo);
7401 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
7411bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7412 SPIRVTypeInst ResType,
7413 MachineInstr &
I)
const {
7416 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7422 Register ExpReg =
I.getOperand(2).getReg();
7424 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7425 SPIRV::OpConvertSToF))
7427 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7434bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7435 SPIRVTypeInst ResType,
7436 MachineInstr &
I)
const {
7452 MachineIRBuilder MIRBuilder(
I);
7453 SPIRVTypeInst FloatType =
7457 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7470 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7471 const bool IsUntyped =
7472 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7474 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7475 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7476 : SPIRV::OpVariable))
7479 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7487 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7490 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7493 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7497 Register IntegralPartReg =
I.getOperand(1).getReg();
7500 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7510 assert(
false &&
"GLSL::Modf is deprecated.");
7521bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7522 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7523 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7524 MachineIRBuilder MIRBuilder(
I);
7525 const SPIRVTypeInst Vec3Ty =
7528 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7540 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7544 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7550 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7557 assert(
I.getOperand(2).isReg());
7558 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7562 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7573bool SPIRVInstructionSelector::loadBuiltinInputID(
7574 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7575 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7576 MachineIRBuilder MIRBuilder(
I);
7578 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7593 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7597 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7606SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7607 MachineInstr &
I)
const {
7608 MachineIRBuilder MIRBuilder(
I);
7619bool SPIRVInstructionSelector::loadHandleBeforePosition(
7620 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7621 MachineInstr &Pos)
const {
7624 Intrinsic::spv_resource_handlefrombinding);
7632 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7633 MachineIRBuilder MIRBuilder(HandleDef);
7634 SPIRVTypeInst VarType = ResType;
7635 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7637 if (IsStructuredBuffer) {
7642 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7644 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7647 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7648 ArraySize, IndexReg, Name, MIRBuilder);
7652 uint32_t LoadOpcode =
7653 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7663bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7664 MachineInstr &
I)
const {
7666 return diagnoseUnsupported(
7667 I,
"this instruction is only supported in shaders.");
7672InstructionSelector *
7676 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
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 void addMemoryOperands(MachineMemOperand *MemOp, MachineInstrBuilder &MIB, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry &GR)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef)
static unsigned getPtrCmpOpcode(unsigned Pred)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
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.
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.
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 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.
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