37#include "llvm/IR/IntrinsicsSPIRV.h"
43#define DEBUG_TYPE "spirv-isel"
50 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
55 std::optional<Register> Bias;
56 std::optional<Register>
Offset;
57 std::optional<Register> MinLod;
58 std::optional<Register> GradX;
59 std::optional<Register> GradY;
60 std::optional<Register> Lod;
61 std::optional<Register> Compare;
64llvm::SPIRV::SelectionControl::SelectionControl
65getSelectionOperandForImm(
int Imm) {
67 return SPIRV::SelectionControl::Flatten;
69 return SPIRV::SelectionControl::DontFlatten;
71 return SPIRV::SelectionControl::None;
75#define GET_GLOBALISEL_PREDICATE_BITSET
76#include "SPIRVGenGlobalISel.inc"
77#undef GET_GLOBALISEL_PREDICATE_BITSET
104#define GET_GLOBALISEL_PREDICATES_DECL
105#include "SPIRVGenGlobalISel.inc"
106#undef GET_GLOBALISEL_PREDICATES_DECL
108#define GET_GLOBALISEL_TEMPORARIES_DECL
109#include "SPIRVGenGlobalISel.inc"
110#undef GET_GLOBALISEL_TEMPORARIES_DECL
134 unsigned BitSetOpcode)
const;
138 unsigned BitSetOpcode)
const;
142 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
149 unsigned Opcode)
const;
152 unsigned Opcode)
const;
173 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
182 Register castPtrToPtrToInt(
Register Ptr, SPIRV::StorageClass::StorageClass SC,
186 bool selectAtomicPtrValue(
202 unsigned OpType)
const;
270 unsigned Opcode)
const;
274 unsigned Opcode)
const;
278 unsigned Opcode)
const;
282 unsigned Opcode)
const;
284 template <
bool Signed>
287 template <
bool Signed>
294 template <
typename PickOpcodeFn>
297 PickOpcodeFn &&PickOpcode)
const;
314 template <
typename PickOpcodeFn>
317 PickOpcodeFn &&PickOpcode)
const;
339 bool IsSigned)
const;
341 bool IsSigned,
unsigned Opcode)
const;
343 bool IsSigned)
const;
349 bool IsSigned)
const;
390 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
391 bool useMISrc =
true,
393 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
394 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
395 bool useMISrc =
true,
397 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
398 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
399 bool setMIFlags =
true,
bool useMISrc =
true,
401 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
402 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
403 bool useMISrc =
true,
406 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
407 MachineInstr &
I)
const;
409 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
412 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
415 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
416 MachineInstr &
I,
unsigned Opcode)
const;
418 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
419 bool WithGroupSync)
const;
421 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
422 MachineInstr &
I)
const;
424 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
425 MachineInstr &
I)
const;
429 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
430 MachineInstr &
I)
const;
432 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
433 MachineInstr &
I)
const;
435 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
436 MachineInstr &
I)
const;
437 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
438 MachineInstr &
I)
const;
439 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
440 SPIRVTypeInst ResType,
441 MachineInstr &
I)
const;
442 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
443 MachineInstr &
I)
const;
446 std::optional<Register> LodReg = std::nullopt)
const;
447 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
450 MachineInstr &
I)
const;
451 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
452 MachineInstr &
I)
const;
453 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
454 MachineInstr &
I)
const;
455 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
458 MachineInstr &
I)
const;
459 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
460 MachineInstr &
I)
const;
461 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
462 SPIRVTypeInst ResType,
463 MachineInstr &
I)
const;
464 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
465 MachineInstr &
I)
const;
466 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
467 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
468 MachineInstr &
I)
const;
469 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
470 MachineInstr &
I)
const;
471 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
472 MachineInstr &
I)
const;
473 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
474 MachineInstr &
I)
const;
475 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
476 MachineInstr &
I)
const;
477 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
478 MachineInstr &
I)
const;
480 bool selectCopySign(
Register ResVReg, SPIRVTypeInst ResType,
481 MachineInstr &
I)
const;
483 bool selectLdexp(
Register ResVReg, SPIRVTypeInst ResType,
484 MachineInstr &
I)
const;
485 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
486 MachineInstr &
I)
const;
487 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
488 MachineInstr &
I)
const;
489 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
490 MachineInstr &
I,
const unsigned DPdOpCode)
const;
492 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
493 SPIRVTypeInst ResType =
nullptr)
const;
495 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
496 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
497 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
499 MachineInstr &
I)
const;
500 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
502 MachineInstr &
I)
const;
504 bool wrapIntoSpecConstantOp(MachineInstr &
I,
507 SPIRVTypeInst getUcharPtrType(MachineInstr &
I,
508 SPIRV::StorageClass::StorageClass SC)
const;
509 Register getUcharPtrTypeReg(MachineInstr &
I,
510 SPIRV::StorageClass::StorageClass SC)
const;
511 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
513 uint32_t Opcode)
const;
514 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
515 SPIRVTypeInst SrcPtrTy)
const;
516 Register buildPointerToResource(SPIRVTypeInst ResType,
517 SPIRV::StorageClass::StorageClass SC,
518 uint32_t Set, uint32_t
Binding,
519 uint32_t ArraySize,
Register IndexReg,
521 MachineIRBuilder MIRBuilder)
const;
522 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
523 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
524 Register &ReadReg, MachineInstr &InsertionPoint)
const;
525 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
528 const ImageOperands *ImOps =
nullptr)
const;
529 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
531 Register CoordinateReg,
const ImageOperands &ImOps,
534 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
535 Register ResVReg, SPIRVTypeInst ResType,
536 MachineInstr &
I)
const;
537 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
538 Register ResVReg, SPIRVTypeInst ResType,
539 MachineInstr &
I)
const;
540 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
541 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
542 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
543 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
546 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
547 Register SrcReg,
unsigned int Opcode,
548 std::function<
bool(
Register, SPIRVTypeInst,
549 MachineInstr &,
Register,
unsigned)>
553bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
555 if (
TET->getTargetExtName() ==
"spirv.Image") {
558 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
559 return TET->getTypeParameter(0)->isIntegerTy();
563#define GET_GLOBALISEL_IMPL
564#include "SPIRVGenGlobalISel.inc"
565#undef GET_GLOBALISEL_IMPL
571 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
574#include
"SPIRVGenGlobalISel.inc"
577#include
"SPIRVGenGlobalISel.inc"
589 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
594 if (HasVRegsReset == &MF)
609 for (
const auto &
MBB : MF) {
610 for (
const auto &
MI :
MBB) {
613 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
620 if (DstType != SrcType)
625 if (DstRC != SrcRC && SrcRC)
637 while (!Stack.empty()) {
642 switch (
MI->getOpcode()) {
643 case TargetOpcode::G_INTRINSIC:
644 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
645 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
648 if (IntrID != Intrinsic::spv_const_composite &&
649 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
653 case TargetOpcode::G_BUILD_VECTOR:
654 case TargetOpcode::G_SPLAT_VECTOR:
656 i < OpDef->getNumOperands(); i++) {
661 Stack.push_back(OpNestedDef);
664 case TargetOpcode::G_CONSTANT:
665 case TargetOpcode::G_FCONSTANT:
666 case TargetOpcode::G_IMPLICIT_DEF:
667 case SPIRV::OpConstantTrue:
668 case SPIRV::OpConstantFalse:
669 case SPIRV::OpConstantI:
670 case SPIRV::OpConstantF:
671 case SPIRV::OpConstantComposite:
672 case SPIRV::OpConstantCompositeContinuedINTEL:
673 case SPIRV::OpConstantSampler:
674 case SPIRV::OpConstantNull:
676 case SPIRV::OpPoisonKHR:
677 case SPIRV::OpConstantFunctionPointerINTEL:
704 case Intrinsic::spv_all:
705 case Intrinsic::spv_alloca:
706 case Intrinsic::spv_any:
707 case Intrinsic::spv_bitcast:
708 case Intrinsic::spv_const_composite:
709 case Intrinsic::spv_degrees:
710 case Intrinsic::spv_distance:
711 case Intrinsic::spv_extractelt:
712 case Intrinsic::spv_extractv:
713 case Intrinsic::spv_faceforward:
714 case Intrinsic::spv_fdot:
715 case Intrinsic::spv_firstbitlow:
716 case Intrinsic::spv_firstbitshigh:
717 case Intrinsic::spv_firstbituhigh:
718 case Intrinsic::spv_frac:
719 case Intrinsic::spv_gep:
720 case Intrinsic::spv_global_offset:
721 case Intrinsic::spv_global_size:
722 case Intrinsic::spv_group_id:
723 case Intrinsic::spv_insertelt:
724 case Intrinsic::spv_insertv:
725 case Intrinsic::spv_isinf:
726 case Intrinsic::spv_isnan:
727 case Intrinsic::spv_isfinite:
728 case Intrinsic::spv_isnormal:
729 case Intrinsic::spv_lerp:
730 case Intrinsic::spv_length:
731 case Intrinsic::spv_normalize:
732 case Intrinsic::spv_num_subgroups:
733 case Intrinsic::spv_num_workgroups:
734 case Intrinsic::spv_ptrcast:
735 case Intrinsic::spv_radians:
736 case Intrinsic::spv_reflect:
737 case Intrinsic::spv_refract:
738 case Intrinsic::spv_resource_getbasepointer:
739 case Intrinsic::spv_resource_getpointer:
740 case Intrinsic::spv_resource_handlefrombinding:
741 case Intrinsic::spv_resource_handlefromimplicitbinding:
742 case Intrinsic::spv_resource_nonuniformindex:
743 case Intrinsic::spv_resource_sample:
744 case Intrinsic::spv_rsqrt:
745 case Intrinsic::spv_saturate:
746 case Intrinsic::spv_sdot:
747 case Intrinsic::spv_sign:
748 case Intrinsic::spv_smoothstep:
749 case Intrinsic::spv_subgroup_id:
750 case Intrinsic::spv_subgroup_local_invocation_id:
751 case Intrinsic::spv_subgroup_max_size:
752 case Intrinsic::spv_subgroup_size:
753 case Intrinsic::spv_thread_id:
754 case Intrinsic::spv_thread_id_in_group:
755 case Intrinsic::spv_udot:
756 case Intrinsic::spv_undef:
757 case Intrinsic::spv_value_md:
758 case Intrinsic::spv_workgroup_size:
770 case SPIRV::OpTypeVoid:
771 case SPIRV::OpTypeBool:
772 case SPIRV::OpTypeInt:
773 case SPIRV::OpTypeFloat:
774 case SPIRV::OpTypeVector:
775 case SPIRV::OpTypeVectorIdEXT:
776 case SPIRV::OpTypeMatrix:
777 case SPIRV::OpTypeImage:
778 case SPIRV::OpTypeSampler:
779 case SPIRV::OpTypeSampledImage:
780 case SPIRV::OpTypeArray:
781 case SPIRV::OpTypeRuntimeArray:
782 case SPIRV::OpTypeStruct:
783 case SPIRV::OpTypeOpaque:
784 case SPIRV::OpTypePointer:
785 case SPIRV::OpTypeFunction:
786 case SPIRV::OpTypeEvent:
787 case SPIRV::OpTypeDeviceEvent:
788 case SPIRV::OpTypeReserveId:
789 case SPIRV::OpTypeQueue:
790 case SPIRV::OpTypePipe:
791 case SPIRV::OpTypeForwardPointer:
792 case SPIRV::OpTypePipeStorage:
793 case SPIRV::OpTypeNamedBarrier:
794 case SPIRV::OpTypeAccelerationStructureNV:
795 case SPIRV::OpTypeCooperativeMatrixNV:
796 case SPIRV::OpTypeCooperativeMatrixKHR:
806 if (
MI.getNumDefs() == 0)
809 for (
const auto &MO :
MI.all_defs()) {
811 if (
Reg.isPhysical()) {
815 for (
const auto &
UseMI :
MRI.use_nodbg_instructions(
Reg)) {
816 if (
UseMI.getOpcode() != SPIRV::OpName) {
823 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
824 MI.isLifetimeMarker()) {
827 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
838 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
839 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
842 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
847 if (
MI.mayStore() ||
MI.isCall() ||
848 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
849 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
850 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
861 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
868void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
870 for (
const auto &MO :
MI.all_defs()) {
874 SmallVector<MachineInstr *, 4> UselessOpNames;
877 "There is still a use of the dead function.");
880 for (MachineInstr *OpNameMI : UselessOpNames) {
882 OpNameMI->eraseFromParent();
887void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
890 removeOpNamesForDeadMI(
MI);
891 MI.eraseFromParent();
894bool SPIRVInstructionSelector::select(MachineInstr &
I) {
895 resetVRegsType(*
I.getParent()->getParent());
897 assert(
I.getParent() &&
"Instruction should be in a basic block!");
898 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
903 removeDeadInstruction(
I);
910 if (Opcode == SPIRV::ASSIGN_TYPE) {
911 Register DstReg =
I.getOperand(0).getReg();
912 Register SrcReg =
I.getOperand(1).getReg();
915 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
916 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
917 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
918 Register SelectDstReg =
Def->getOperand(0).getReg();
919 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
921 assert(SuccessToSelectSelect);
923 Def->eraseFromParent();
930 bool Res = selectImpl(
I, *CoverageInfo);
932 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
933 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
937 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
949 }
else if (
I.getNumDefs() == 1) {
961 removeDeadInstruction(
I);
966 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
967 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
973 bool HasDefs =
I.getNumDefs() > 0;
976 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
977 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
978 if (spvSelect(ResVReg, ResType,
I)) {
980 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
991 case TargetOpcode::G_CONSTANT:
992 case TargetOpcode::G_FCONSTANT:
999 MachineInstr &
I)
const {
1002 if (DstRC != SrcRC && SrcRC)
1004 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
1011bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
1012 SPIRVTypeInst ResType,
1013 MachineInstr &
I)
const {
1014 const unsigned Opcode =
I.getOpcode();
1016 return selectImpl(
I, *CoverageInfo);
1018 case TargetOpcode::G_CONSTANT:
1019 case TargetOpcode::G_FCONSTANT:
1020 return selectConst(ResVReg, ResType,
I);
1021 case TargetOpcode::G_GLOBAL_VALUE:
1022 return selectGlobalValue(ResVReg,
I);
1023 case TargetOpcode::G_IMPLICIT_DEF:
1024 return selectOpUndef(ResVReg, ResType,
I);
1025 case TargetOpcode::G_FREEZE:
1026 return selectFreeze(ResVReg, ResType,
I);
1028 case TargetOpcode::G_INTRINSIC:
1029 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1030 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1031 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1032 return selectIntrinsic(ResVReg, ResType,
I);
1033 case TargetOpcode::G_BITREVERSE:
1034 return selectBitreverse(ResVReg, ResType,
I);
1036 case TargetOpcode::G_BUILD_VECTOR:
1037 return selectBuildVector(ResVReg, ResType,
I);
1038 case TargetOpcode::G_SPLAT_VECTOR:
1039 return selectSplatVector(ResVReg, ResType,
I);
1040 case TargetOpcode::G_CONCAT_VECTORS:
1041 return selectConcatVectors(ResVReg, ResType,
I);
1043 case TargetOpcode::G_SHUFFLE_VECTOR: {
1044 MachineBasicBlock &BB = *
I.getParent();
1045 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1048 .
addUse(
I.getOperand(1).getReg())
1049 .
addUse(
I.getOperand(2).getReg());
1050 for (
auto V :
I.getOperand(3).getShuffleMask())
1055 case TargetOpcode::G_MEMMOVE:
1056 case TargetOpcode::G_MEMCPY:
1057 case TargetOpcode::G_MEMCPY_INLINE:
1058 case TargetOpcode::G_MEMSET:
1059 case TargetOpcode::G_MEMSET_INLINE:
1060 return selectMemOperation(ResVReg,
I);
1062 case TargetOpcode::G_ICMP:
1063 return selectICmp(ResVReg, ResType,
I);
1064 case TargetOpcode::G_FCMP:
1065 return selectFCmp(ResVReg, ResType,
I);
1067 case TargetOpcode::G_FRAME_INDEX:
1068 return selectFrameIndex(ResVReg, ResType,
I);
1070 case TargetOpcode::G_LOAD:
1071 return selectLoad(ResVReg, ResType,
I);
1072 case TargetOpcode::G_STORE:
1073 return selectStore(
I);
1075 case TargetOpcode::G_BR:
1076 return selectBranch(
I);
1077 case TargetOpcode::G_BRCOND:
1078 return selectBranchCond(
I);
1080 case TargetOpcode::G_PHI:
1081 return selectPhi(ResVReg,
I);
1083 case TargetOpcode::G_FPTOSI:
1084 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1085 case TargetOpcode::G_FPTOUI:
1086 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1088 case TargetOpcode::G_FPTOSI_SAT:
1089 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1090 case TargetOpcode::G_FPTOUI_SAT:
1091 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1093 case TargetOpcode::G_SITOFP:
1094 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1095 case TargetOpcode::G_UITOFP:
1096 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1098 case TargetOpcode::G_CTPOP:
1099 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1100 case TargetOpcode::G_SMIN:
1101 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1102 case TargetOpcode::G_UMIN:
1103 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1105 case TargetOpcode::G_SMAX:
1106 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1107 case TargetOpcode::G_UMAX:
1108 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1110 case TargetOpcode::G_SCMP:
1111 return selectSUCmp(ResVReg, ResType,
I,
true);
1112 case TargetOpcode::G_UCMP:
1113 return selectSUCmp(ResVReg, ResType,
I,
false);
1114 case TargetOpcode::G_LROUND:
1115 case TargetOpcode::G_LLROUND: {
1120 regForLround, *(
I.getParent()->getParent()));
1122 CL::round, GL::Round,
false);
1124 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1131 case TargetOpcode::G_STRICT_FMA:
1132 case TargetOpcode::G_FMA: {
1135 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1138 .
addUse(
I.getOperand(1).getReg())
1139 .
addUse(
I.getOperand(2).getReg())
1140 .
addUse(
I.getOperand(3).getReg())
1145 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1148 case TargetOpcode::G_FLDEXP:
1149 case TargetOpcode::G_STRICT_FLDEXP:
1150 return selectLdexp(ResVReg, ResType,
I);
1152 case TargetOpcode::G_FPOW:
1153 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1154 case TargetOpcode::G_FPOWI:
1155 return selectFpowi(ResVReg, ResType,
I);
1157 case TargetOpcode::G_FEXP:
1158 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1159 case TargetOpcode::G_FEXP2:
1160 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1161 case TargetOpcode::G_FEXP10:
1162 return selectExp10(ResVReg, ResType,
I);
1164 case TargetOpcode::G_FMODF:
1165 return selectModf(ResVReg, ResType,
I);
1166 case TargetOpcode::G_FSINCOS:
1167 return selectSincos(ResVReg, ResType,
I);
1169 case TargetOpcode::G_FLOG:
1170 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1171 case TargetOpcode::G_FLOG2:
1172 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1173 case TargetOpcode::G_FLOG10:
1174 return selectLog10(ResVReg, ResType,
I);
1176 case TargetOpcode::G_FABS:
1177 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1178 case TargetOpcode::G_ABS:
1179 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1181 case TargetOpcode::G_FMINNUM:
1182 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1183 case TargetOpcode::G_FMAXNUM:
1184 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1186 case TargetOpcode::G_FCOPYSIGN:
1187 return selectCopySign(ResVReg, ResType,
I);
1189 case TargetOpcode::G_FCEIL:
1190 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1191 case TargetOpcode::G_FFLOOR:
1192 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1194 case TargetOpcode::G_FCOS:
1195 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1196 case TargetOpcode::G_FSIN:
1197 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1198 case TargetOpcode::G_FTAN:
1199 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1200 case TargetOpcode::G_FACOS:
1201 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1202 case TargetOpcode::G_FASIN:
1203 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1204 case TargetOpcode::G_FATAN:
1205 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1206 case TargetOpcode::G_FATAN2:
1207 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1208 case TargetOpcode::G_FCOSH:
1209 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1210 case TargetOpcode::G_FSINH:
1211 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1212 case TargetOpcode::G_FTANH:
1213 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1215 case TargetOpcode::G_STRICT_FSQRT:
1216 case TargetOpcode::G_FSQRT:
1217 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1219 case TargetOpcode::G_CTTZ:
1220 case TargetOpcode::G_CTTZ_ZERO_POISON:
1221 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1222 case TargetOpcode::G_CTLZ:
1223 case TargetOpcode::G_CTLZ_ZERO_POISON:
1224 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1226 case TargetOpcode::G_INTRINSIC_ROUND:
1227 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1228 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1229 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1230 case TargetOpcode::G_INTRINSIC_TRUNC:
1231 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1232 case TargetOpcode::G_FRINT:
1233 case TargetOpcode::G_FNEARBYINT:
1234 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1236 case TargetOpcode::G_SMULH:
1237 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1238 case TargetOpcode::G_UMULH:
1239 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1241 case TargetOpcode::G_SADDSAT:
1242 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1243 case TargetOpcode::G_UADDSAT:
1244 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1245 case TargetOpcode::G_SSUBSAT:
1246 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1247 case TargetOpcode::G_USUBSAT:
1248 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1250 case TargetOpcode::G_FFREXP:
1251 return selectFrexp(ResVReg, ResType,
I);
1253 case TargetOpcode::G_UADDO:
1254 return selectOverflowArith(ResVReg, ResType,
I,
1256 : SPIRV::OpIAddCarryS);
1257 case TargetOpcode::G_USUBO:
1258 return selectOverflowArith(ResVReg, ResType,
I,
1260 : SPIRV::OpISubBorrowS);
1261 case TargetOpcode::G_UMULO:
1262 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1263 case TargetOpcode::G_SMULO:
1264 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1266 case TargetOpcode::G_SEXT:
1267 return selectExt(ResVReg, ResType,
I,
true);
1268 case TargetOpcode::G_ANYEXT:
1269 case TargetOpcode::G_ZEXT:
1270 return selectExt(ResVReg, ResType,
I,
false);
1271 case TargetOpcode::G_TRUNC:
1272 return selectTrunc(ResVReg, ResType,
I);
1273 case TargetOpcode::G_FPTRUNC:
1274 case TargetOpcode::G_FPEXT:
1275 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1277 case TargetOpcode::G_PTRTOINT:
1278 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1279 case TargetOpcode::G_INTTOPTR:
1280 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1281 case TargetOpcode::G_BITCAST:
1282 return selectBitcast(ResVReg, ResType,
I);
1283 case TargetOpcode::G_ADDRSPACE_CAST:
1284 return selectAddrSpaceCast(ResVReg, ResType,
I);
1285 case TargetOpcode::G_PTRMASK:
1286 return selectPtrMask(ResVReg, ResType,
I);
1287 case TargetOpcode::G_PTR_ADD: {
1289 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1293 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1294 (*II).getOpcode() == TargetOpcode::COPY ||
1295 (*II).getOpcode() == SPIRV::OpVariable ||
1296 (*II).getOpcode() == SPIRV::OpUntypedVariableKHR) &&
1299 bool IsGVInit =
false;
1303 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1304 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1305 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1306 (*UseIt).getOpcode() == SPIRV::OpVariable ||
1307 (*UseIt).getOpcode() == SPIRV::OpUntypedVariableKHR) {
1319 const bool UseUntypedPointers =
1320 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1321 if (UseUntypedPointers) {
1322 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1325 .
addImm(
static_cast<uint32_t
>(
1326 SPIRV::Opcode::UntypedInBoundsPtrAccessChainKHR))
1329 .
addUse(
I.getOperand(2).getReg())
1336 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1348 return diagnoseUnsupported(
1349 I,
"incompatible result and operand types in a bitcast");
1351 MachineInstrBuilder MIB =
1352 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1359 : SPIRV::OpInBoundsPtrAccessChain))
1363 .
addUse(
I.getOperand(2).getReg())
1366 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1370 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1372 .
addUse(
I.getOperand(2).getReg())
1381 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1384 .
addImm(
static_cast<uint32_t
>(
1385 SPIRV::Opcode::InBoundsPtrAccessChain))
1388 .
addUse(
I.getOperand(2).getReg());
1393 case TargetOpcode::G_ATOMICRMW_OR:
1394 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1395 case TargetOpcode::G_ATOMICRMW_ADD:
1396 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1397 case TargetOpcode::G_ATOMICRMW_AND:
1398 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1399 case TargetOpcode::G_ATOMICRMW_MAX:
1400 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1401 case TargetOpcode::G_ATOMICRMW_MIN:
1402 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1403 case TargetOpcode::G_ATOMICRMW_SUB:
1404 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1405 case TargetOpcode::G_ATOMICRMW_XOR:
1406 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1407 case TargetOpcode::G_ATOMICRMW_UMAX:
1408 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1409 case TargetOpcode::G_ATOMICRMW_UMIN:
1410 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1411 case TargetOpcode::G_ATOMICRMW_XCHG:
1412 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1414 case TargetOpcode::G_ATOMICRMW_FADD:
1415 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1416 case TargetOpcode::G_ATOMICRMW_FSUB:
1418 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1420 : SPIRV::OpFNegate);
1421 case TargetOpcode::G_ATOMICRMW_FMIN:
1422 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1423 case TargetOpcode::G_ATOMICRMW_FMAX:
1424 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1426 case TargetOpcode::G_FENCE:
1427 return selectFence(
I);
1429 case TargetOpcode::G_STACKSAVE:
1430 return selectStackSave(ResVReg, ResType,
I);
1431 case TargetOpcode::G_STACKRESTORE:
1432 return selectStackRestore(
I);
1434 case TargetOpcode::G_UNMERGE_VALUES:
1437 case TargetOpcode::G_TRAP:
1438 case TargetOpcode::G_UBSANTRAP:
1439 return selectTrap(
I);
1444 case TargetOpcode::DBG_LABEL:
1446 case TargetOpcode::G_DEBUGTRAP:
1447 return selectDebugTrap(ResVReg, ResType,
I);
1448 case TargetOpcode::G_PREFETCH:
1449 return selectPrefetch(
I);
1456bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1457 SPIRVTypeInst ResType,
1458 MachineInstr &
I)
const {
1459 unsigned Opcode = SPIRV::OpNop;
1466bool SPIRVInstructionSelector::selectPrefetch(MachineInstr &
I)
const {
1475 MachineIRBuilder MIRBuilder(
I);
1477 const SPIRVTypeInst PointerSizeType =
1485 Register AddrVal =
I.getOperand(0).getReg();
1488 return selectExtInst(ExtReg, GR.
getOpTypeVoid(MIRBuilder),
I, CL::prefetch,
1490 {AddrVal, ConstIntOne});
1495bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1496 SPIRVTypeInst ResType,
1498 GL::GLSLExtInst GLInst,
1499 bool setMIFlags,
bool useMISrc,
1502 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1503 return diagnoseUnsupported(
1505 "this instruction is only supported with the GLSL extended instruction "
1507 return selectExtInst(ResVReg, ResType,
I,
1508 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1509 setMIFlags, useMISrc, SrcRegs);
1512bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1513 SPIRVTypeInst ResType,
1515 CL::OpenCLExtInst CLInst,
1516 bool setMIFlags,
bool useMISrc,
1518 return selectExtInst(ResVReg, ResType,
I,
1519 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1520 setMIFlags, useMISrc, SrcRegs);
1523bool SPIRVInstructionSelector::selectExtInst(
1524 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1525 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1527 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1528 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1529 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1533bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1534 SPIRVTypeInst ResType,
1537 bool setMIFlags,
bool useMISrc,
1540 for (
const auto &[InstructionSet, Opcode] : Insts) {
1544 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1547 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1552 const unsigned NumOps =
I.getNumOperands();
1555 I.getOperand(Index).getType() ==
1556 MachineOperand::MachineOperandType::MO_IntrinsicID)
1559 MIB.
add(
I.getOperand(Index));
1571bool SPIRVInstructionSelector::selectCopySign(
Register ResVReg,
1572 SPIRVTypeInst ResType,
1573 MachineInstr &
I)
const {
1575 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1580 Register MagnitudeReg =
I.getOperand(1).getReg();
1581 Register SignReg =
I.getOperand(2).getReg();
1589 unsigned AndOpcode, OrOpcode;
1590 if (ComponentCount > 1) {
1594 AndOpcode = SPIRV::OpBitwiseAndV;
1595 OrOpcode = SPIRV::OpBitwiseOrV;
1599 AndOpcode = SPIRV::OpBitwiseAndS;
1600 OrOpcode = SPIRV::OpBitwiseOrS;
1606 return selectOpWithSrcs(ResReg, IntType,
I, SrcRegs, Opcode);
1609 Register MagnitudeInt, SignInt, MagnitudeBits, SignBits, CombinedInt;
1610 if (!EmitBitOp(MagnitudeInt, {MagnitudeReg}, SPIRV::OpBitcast) ||
1611 !EmitBitOp(SignInt, {SignReg}, SPIRV::OpBitcast) ||
1612 !EmitBitOp(MagnitudeBits, {MagnitudeInt, NotSignMask}, AndOpcode) ||
1613 !EmitBitOp(SignBits, {SignInt, SignMask}, AndOpcode) ||
1614 !EmitBitOp(CombinedInt, {MagnitudeBits, SignBits}, OrOpcode))
1617 return selectOpWithSrcs(ResVReg, ResType,
I, {CombinedInt}, SPIRV::OpBitcast);
1620bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1621 SPIRVTypeInst ResType,
1622 MachineInstr &
I)
const {
1623 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1624 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1625 for (
const auto &Ex : ExtInsts) {
1626 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1627 uint32_t Opcode = Ex.second;
1631 MachineIRBuilder MIRBuilder(
I);
1638 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1645 const bool IsUntyped =
1646 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1648 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1649 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1650 : SPIRV::OpVariable))
1653 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1658 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1661 .
addImm(
static_cast<uint32_t
>(Ex.first))
1663 .
add(
I.getOperand(2))
1667 Register ExpResReg =
I.getOperand(1).getReg();
1669 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1679bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1680 SPIRVTypeInst ResType,
1681 MachineInstr &
I)
const {
1682 Register XReg =
I.getOperand(1).getReg();
1683 Register ExpReg =
I.getOperand(2).getReg();
1689 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1690 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1692 ExpReg = buildVectorSplat(ExpReg, NumElts,
I);
1695 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1696 true,
false, {XReg, ExpReg});
1699bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1700 SPIRVTypeInst ResType,
1701 MachineInstr &
I)
const {
1702 Register CosResVReg =
I.getOperand(1).getReg();
1703 unsigned SrcIdx =
I.getNumExplicitDefs();
1708 MachineIRBuilder MIRBuilder(
I);
1710 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1717 const bool IsUntyped =
1718 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1720 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1721 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1722 : SPIRV::OpVariable))
1725 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1729 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1732 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1734 .
add(
I.getOperand(SrcIdx))
1738 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1746 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1749 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1751 .
add(
I.getOperand(SrcIdx))
1753 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1756 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1758 .
add(
I.getOperand(SrcIdx))
1765bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1766 SPIRVTypeInst ResType,
1769 unsigned Opcode)
const {
1770 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1780bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1781 SPIRVTypeInst ResType,
1784 unsigned Opcode)
const {
1785 MachineIRBuilder MIRBuilder(
I);
1787 Register OpReg =
I.getOperand(1).getReg();
1793 SPIRVTypeInst ExtType =
1800 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1804 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1807 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1810bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1811 SPIRVTypeInst ResType,
1814 unsigned Opcode)
const {
1815 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1818bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1819 SPIRVTypeInst ResType,
1822 unsigned Opcode)
const {
1823 MachineIRBuilder MIRBuilder(
I);
1832 SPIRVTypeInst WorkingType =
1839 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {SrcReg}, SPIRV::OpUConvert))
1843 if (!selectOpWithSrcs(LowCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1851 IsScalar ? SPIRV::OpShiftRightLogicalS : SPIRV::OpShiftRightLogicalV;
1853 if (!selectOpWithSrcs(Shift, SrcType,
I, {SrcReg, ShiftAmount}, ShiftOp))
1857 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {Shift}, SPIRV::OpUConvert))
1861 if (!selectOpWithSrcs(HighCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1866 if (!selectOpWithSrcs(Sum, WorkingType,
I, {HighCount, LowCount},
1867 IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV))
1871 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1872 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1875bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1876 SPIRVTypeInst ResType,
1878 unsigned Opcode)
const {
1883 if (!STI.getTargetTriple().isVulkanOS())
1884 return selectUnOp(ResVReg, ResType,
I, Opcode);
1886 Register OpReg =
I.getOperand(1).getReg();
1889 : SPIRV::OpUConvert;
1893 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1895 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1897 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1899 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1903bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1904 SPIRVTypeInst ResType,
1906 unsigned Opcode)
const {
1908 Register SrcReg =
I.getOperand(1).getReg();
1913 unsigned DefOpCode = DefIt->getOpcode();
1914 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1917 if (
auto *VRD =
getVRegDef(*
MRI, DefIt->getOperand(1).getReg()))
1918 DefOpCode = VRD->getOpcode();
1920 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1921 DefOpCode == TargetOpcode::G_CONSTANT ||
1922 DefOpCode == SPIRV::OpVariable ||
1923 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1924 DefOpCode == SPIRV::OpConstantI) {
1930 uint32_t SpecOpcode = 0;
1932 case SPIRV::OpConvertPtrToU:
1933 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1935 case SPIRV::OpConvertUToPtr:
1936 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1941 TII.get(SPIRV::OpSpecConstantOp))
1951 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1955bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1956 SPIRVTypeInst ResType,
1957 MachineInstr &
I)
const {
1958 Register OpReg =
I.getOperand(1).getReg();
1959 SPIRVTypeInst OpType =
1962 return diagnoseUnsupported(
1963 I,
"incompatible result and operand types in a bitcast");
1964 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1970 std::optional<Align> AlignOverride = std::nullopt) {
1975 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1976 if (
MemOp->isNonTemporal())
1977 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1979 if (!ST->isShader() &&
MemOp->getAlign().value())
1980 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1984 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1985 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1989 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1991 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1995 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1999 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
2001 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
2002 MIB.
addImm(AlignOverride.value_or(
MemOp->getAlign()).value());
2013 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2015 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2017 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2021bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2022 SPIRVTypeInst ResType,
2023 MachineInstr &
I)
const {
2025 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2030 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2031 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2033 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2035 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2039 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2043 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2044 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2045 I.getDebugLoc(),
I);
2049 MachineIRBuilder MIRBuilder(
I);
2051 if (
I.getNumMemOperands()) {
2052 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2053 if (MemOp->isAtomic())
2054 return selectAtomicLoad(ResVReg, ResType,
I);
2057 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2061 if (!
I.getNumMemOperands()) {
2062 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2064 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2073Register SPIRVInstructionSelector::createPtrSizedIntReg(
2074 MachineIRBuilder &MIRBuilder)
const {
2075 SPIRVTypeInst IntType =
2085SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2086 MachineIRBuilder &MIRBuilder)
const {
2087 SPIRVTypeInst IntType =
2089 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2090 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2098Register SPIRVInstructionSelector::castPtrToPtrToInt(
2099 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2100 MachineIRBuilder &MIRBuilder)
const {
2101 SPIRVTypeInst IntType =
2103 SPIRVTypeInst PtrType =
2117bool SPIRVInstructionSelector::selectAtomicPtrValue(
2118 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2119 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2127 Register IntResult = EmitAtomic(IntType);
2129 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2137bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2138 SPIRVTypeInst ResType,
2139 MachineInstr &
I)
const {
2140 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2143 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2146 return diagnoseUnsupported(
2147 I,
"Lowering to SPIR-V of atomic load is only "
2148 "allowed for integer, floating point or pointer types");
2150 assert(
I.getNumMemOperands());
2151 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2152 assert(MemOp.isAtomic());
2154 uint32_t
Scope =
static_cast<uint32_t
>(
2156 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2162 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2163 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2166 Register MemSemReg = buildI32ConstantInEntryBlock(Sem,
I);
2168 MachineIRBuilder MIRBuilder(
I);
2172 return diagnoseUnsupported(
2173 I,
"Lowering to SPIR-V of atomic load is only "
2174 "allowed for pointer types for physical addressing model");
2179 SPIRV::StorageClass::StorageClass SC =
2181 return selectAtomicPtrValue(
2182 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2183 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2184 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2195 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2206bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2208 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2209 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2214 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2215 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2217 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2222 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2226 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2227 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2228 SPIRVTypeInst SampledType =
2230 SPIRVTypeInst StoreValCompType =
2232 if (StoreValCompType && StoreValCompType != SampledType) {
2235 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2238 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2243 StoreVal = PackedReg;
2246 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2247 TII.get(SPIRV::OpImageWrite))
2253 if (sampledTypeIsSignedInteger(LLVMHandleType))
2256 BMI.constrainAllUses(
TII,
TRI, RBI);
2263 if (PointeeTy && PointeeTy->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
2264 StoreTy->
getOpcode() != SPIRV::OpTypeVectorIdEXT &&
2278 if (
I.getNumMemOperands()) {
2279 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2280 if (MemOp->isAtomic())
2281 return selectAtomicStore(
I);
2288 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2289 PtrSC == SPIRV::StorageClass::Input ||
2290 PtrSC == SPIRV::StorageClass::PushConstant)
2291 return diagnoseUnsupported(
2292 I,
"store into a read-only SPIR-V storage class is not allowed");
2294 MachineIRBuilder MIRBuilder(
I);
2296 if (!
I.getNumMemOperands()) {
2297 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2299 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2308bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2309 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2312 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2313 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2318 if (!PointeeType && PtrType &&
2319 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2322 return diagnoseUnsupported(
I,
2323 "Lowering to SPIR-V of atomic store is only "
2324 "allowed for integer or floating point types");
2326 assert(
I.getNumMemOperands());
2327 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2328 assert(MemOp.isAtomic());
2330 uint32_t
Scope =
static_cast<uint32_t
>(
2332 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2338 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2339 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2342 Register MemSemReg = buildI32ConstantInEntryBlock(Sem,
I);
2343 MachineIRBuilder MIRBuilder(
I);
2347 return diagnoseUnsupported(
2348 I,
"Lowering to SPIR-V of atomic store is only "
2349 "allowed for pointer types for physical addressing model");
2354 SPIRV::StorageClass::StorageClass SC =
2356 return selectAtomicPtrValue(
2357 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2359 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2372 return diagnoseUnsupported(
I,
2373 "Lowering to SPIR-V of atomic store is only "
2374 "allowed for integer or floating point types");
2376 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2386bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2387 SPIRVTypeInst ResType,
2388 MachineInstr &
I)
const {
2389 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2397 const Register PtrsReg =
I.getOperand(2).getReg();
2398 const uint32_t
Alignment =
I.getOperand(3).getImm();
2399 const Register MaskReg =
I.getOperand(4).getReg();
2400 const Register PassthruReg =
I.getOperand(5).getReg();
2401 const Register AlignmentReg = buildI32ConstantInEntryBlock(Alignment,
I);
2405 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2416bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2417 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2424 const Register ValuesReg =
I.getOperand(1).getReg();
2425 const Register PtrsReg =
I.getOperand(2).getReg();
2426 const uint32_t
Alignment =
I.getOperand(3).getImm();
2427 const Register MaskReg =
I.getOperand(4).getReg();
2428 const Register AlignmentReg = buildI32ConstantInEntryBlock(Alignment,
I);
2432 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2441bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2442 const Twine &
Msg)
const {
2443 const Function &
F =
I.getMF()->getFunction();
2444 F.getContext().diagnose(
2445 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2449bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2450 SPIRVTypeInst ResType,
2451 MachineInstr &
I)
const {
2452 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2453 return diagnoseUnsupported(
2454 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2455 "SPIR-V extension: SPV_INTEL_variable_length_array");
2457 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2464bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2465 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2466 return diagnoseUnsupported(
2468 "llvm.stackrestore intrinsic: this instruction requires the following "
2469 "SPIR-V extension: SPV_INTEL_variable_length_array");
2470 if (!
I.getOperand(0).isReg())
2473 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2474 .
addUse(
I.getOperand(0).getReg())
2480SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2481 MachineIRBuilder MIRBuilder(
I);
2482 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2489 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2493 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2494 Type *ArrTy = ArrayType::get(ValTy, Num);
2496 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2499 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2510 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2511 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2512 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2513 : SPIRV::OpVariable))
2516 .
addImm(SPIRV::StorageClass::UniformConstant);
2529bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2532 Register DstReg =
I.getOperand(0).getReg();
2538 return diagnoseUnsupported(
2539 I,
"OpCopyMemory requires operands to have the same type");
2544 return diagnoseUnsupported(
2545 I,
"Unable to determine pointee type size for OpCopyMemory");
2546 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2547 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2548 return diagnoseUnsupported(
2549 I,
"OpCopyMemory requires the size to match the pointee type size");
2552 const unsigned Opcode =
2553 IsLogical ? SPIRV::OpCopyMemory : SPIRV::OpCopyMemorySized;
2554 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2558 MIB.
addUse(
I.getOperand(2).getReg());
2560 if (
I.getNumMemOperands()) {
2561 MachineIRBuilder MIRBuilder(
I);
2562 const MachineMemOperand *DstMemOp = *
I.memoperands_begin();
2564 Align SrcAlign = DstAlign;
2567 if (
I.getNumMemOperands() > 1 && !STI.
isShader())
2568 SrcAlign = (*std::next(
I.memoperands_begin()))->
getAlign();
2578 std::min(DstAlign, SrcAlign));
2585bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2586 MachineInstr &
I)
const {
2588 Register SizeReg =
I.getOperand(2).getReg();
2590 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2594 Register SrcReg =
I.getOperand(1).getReg();
2595 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2596 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2597 Register VarReg = getOrCreateMemSetGlobal(
I);
2600 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2602 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2604 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2607 if (!selectCopyMemory(
I, SrcReg))
2609 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2610 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2615bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2616 SPIRVTypeInst ResType,
2619 unsigned NegateOpcode)
const {
2621 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2622 uint32_t
Scope =
static_cast<uint32_t
>(
2624 MemOp->getSyncScopeID()));
2625 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2627 Register Ptr =
I.getOperand(1).getReg();
2628 uint32_t ScSem =
static_cast<uint32_t
>(
2632 Register MemSemReg = buildI32ConstantInEntryBlock(
2636 Register ValueReg =
I.getOperand(2).getReg();
2637 if (NegateOpcode != 0) {
2640 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2646 if (NewOpcode != SPIRV::OpAtomicExchange)
2647 return diagnoseUnsupported(
2648 I,
"Lowering to SPIR-V of this atomic operation is not "
2649 "allowed for pointer types");
2651 return diagnoseUnsupported(
2652 I,
"Lowering to SPIR-V of atomic exchange is only "
2653 "allowed for pointer types for physical addressing model");
2660 MachineIRBuilder MIRBuilder(
I);
2662 return selectAtomicPtrValue(
2663 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2665 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2666 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2667 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2675 return ExchangeResReg;
2679 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2690bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2691 unsigned ArgI =
I.getNumOperands() - 1;
2693 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2694 SPIRVTypeInst SrcType =
2698 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2702 unsigned CurrentIndex = 0;
2703 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2704 Register ResVReg =
I.getOperand(i).getReg();
2713 ResType = ScalarType;
2722 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2728 for (
unsigned j = 0;
j < NumElements; ++
j) {
2729 MIB.
addImm(CurrentIndex + j);
2731 CurrentIndex += NumElements;
2735 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2747bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2750 ? SPIRV::MemorySemantics::UniformMemory |
2751 SPIRV::MemorySemantics::WorkgroupMemory |
2752 SPIRV::MemorySemantics::ImageMemory
2753 : SPIRV::MemorySemantics::WorkgroupMemory |
2754 SPIRV::MemorySemantics::CrossWorkgroupMemory |
2755 SPIRV::MemorySemantics::ImageMemory;
2757 STI.getTargetTriple(),
static_cast<uint32_t
>(
getMemSemantics(AO)), ScSem);
2758 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
2762 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2764 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2771bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2772 SPIRVTypeInst ResType,
2774 unsigned Opcode)
const {
2775 Type *ResTy =
nullptr;
2778 return diagnoseUnsupported(
2780 "Not enough info to select the arithmetic with overflow instruction");
2782 return diagnoseUnsupported(
I,
2783 "Expect struct type result for the arithmetic "
2784 "with overflow instruction");
2790 MachineIRBuilder MIRBuilder(
I);
2792 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2793 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2800 Register ZeroReg = buildZerosVal(ResType,
I);
2805 if (ResName.
size() > 0)
2813 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2814 MIB.
addUse(
I.getOperand(i).getReg());
2820 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2822 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2823 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2830 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2831 .
addDef(
I.getOperand(1).getReg())
2839bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2840 SPIRVTypeInst ResType,
2841 MachineInstr &
I)
const {
2843 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2844 Register Ptr =
I.getOperand(2).getReg();
2845 Register ScopeReg =
I.getOperand(5).getReg();
2846 Register MemSemEqReg =
I.getOperand(6).getReg();
2847 Register MemSemNeqReg =
I.getOperand(7).getReg();
2849 Register Val =
I.getOperand(4).getReg();
2853 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2872 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2879 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2891 case SPIRV::StorageClass::DeviceOnlyINTEL:
2892 case SPIRV::StorageClass::HostOnlyINTEL:
2901 bool IsGRef =
false;
2902 bool IsAllowedRefs =
2903 llvm::all_of(
MRI->use_instructions(ResVReg), [&IsGRef](
auto const &It) {
2904 unsigned Opcode = It.getOpcode();
2905 if (Opcode == SPIRV::OpConstantComposite ||
2906 Opcode == SPIRV::OpSpecConstantComposite ||
2907 Opcode == SPIRV::OpVariable ||
2908 Opcode == SPIRV::OpUntypedVariableKHR ||
2909 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2910 return IsGRef = true;
2911 return Opcode == SPIRV::OpName;
2913 return IsAllowedRefs && IsGRef;
2916SPIRVTypeInst SPIRVInstructionSelector::getUcharPtrType(
2917 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2919 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC);
2922Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2923 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2928SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2930 uint32_t Opcode)
const {
2931 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2932 TII.get(SPIRV::OpSpecConstantOp))
2940SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2941 SPIRVTypeInst SrcPtrTy)
const {
2942 SPIRVTypeInst GenericPtrTy =
2946 SPIRV::StorageClass::Generic),
2950 MachineInstrBuilder MIB = buildSpecConstantOp(
2952 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2962bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2963 SPIRVTypeInst ResType,
2964 MachineInstr &
I)
const {
2968 Register SrcPtr =
I.getOperand(1).getReg();
2973 return BuildCOPY(ResVReg, SrcPtr,
I);
2983 unsigned SpecOpcode = [&]() ->
unsigned {
2984 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2985 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2987 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2989 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
3000 if (SrcSC != SPIRV::StorageClass::CodeSectionINTEL &&
3001 (SpecOpcode || ViaGeneric)) {
3002 SPIRVTypeInst UcharSrcPtrTy = getUcharPtrType(
I, SrcSC);
3003 if (SrcPtrTy != UcharSrcPtrTy) {
3006 MachineInstrBuilder MIB = buildSpecConstantOp(
3008 static_cast<uint32_t
>(SPIRV::Opcode::Bitcast));
3012 SrcPtrTy = UcharSrcPtrTy;
3016 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
3018 .constrainAllUses(
TII,
TRI, RBI);
3019 }
else if (ViaGeneric) {
3020 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
3022 buildSpecConstantOp(
3024 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
3025 .constrainAllUses(
TII,
TRI, RBI);
3032 return BuildCOPY(ResVReg, SrcPtr,
I);
3034 if ((SrcSC == SPIRV::StorageClass::Function &&
3035 DstSC == SPIRV::StorageClass::Private) ||
3036 (DstSC == SPIRV::StorageClass::Function &&
3037 SrcSC == SPIRV::StorageClass::Private))
3038 return BuildCOPY(ResVReg, SrcPtr,
I);
3042 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3045 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3048 SPIRVTypeInst GenericPtrTy =
3067 return selectUnOp(ResVReg, ResType,
I,
3068 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
3070 return selectUnOp(ResVReg, ResType,
I,
3071 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
3073 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3075 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3085bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3086 SPIRVTypeInst ResType,
3087 MachineInstr &
I)
const {
3089 return diagnoseUnsupported(
3090 I,
"G_PTRMASK is not supported with logical SPIR-V");
3095 Register PtrReg =
I.getOperand(1).getReg();
3096 Register MaskReg =
I.getOperand(2).getReg();
3115 ? SPIRV::OpBitwiseAndV
3116 : SPIRV::OpBitwiseAndS;
3139 return SPIRV::OpFOrdEqual;
3141 return SPIRV::OpFOrdGreaterThanEqual;
3143 return SPIRV::OpFOrdGreaterThan;
3145 return SPIRV::OpFOrdLessThanEqual;
3147 return SPIRV::OpFOrdLessThan;
3149 return SPIRV::OpFOrdNotEqual;
3151 return SPIRV::OpOrdered;
3153 return SPIRV::OpFUnordEqual;
3155 return SPIRV::OpFUnordGreaterThanEqual;
3157 return SPIRV::OpFUnordGreaterThan;
3159 return SPIRV::OpFUnordLessThanEqual;
3161 return SPIRV::OpFUnordLessThan;
3163 return SPIRV::OpFUnordNotEqual;
3165 return SPIRV::OpUnordered;
3175 return SPIRV::OpIEqual;
3177 return SPIRV::OpINotEqual;
3179 return SPIRV::OpSGreaterThanEqual;
3181 return SPIRV::OpSGreaterThan;
3183 return SPIRV::OpSLessThanEqual;
3185 return SPIRV::OpSLessThan;
3187 return SPIRV::OpUGreaterThanEqual;
3189 return SPIRV::OpUGreaterThan;
3191 return SPIRV::OpULessThanEqual;
3193 return SPIRV::OpULessThan;
3202 return SPIRV::OpPtrEqual;
3204 return SPIRV::OpPtrNotEqual;
3215 return SPIRV::OpLogicalEqual;
3217 return SPIRV::OpLogicalNotEqual;
3255bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3256 SPIRVTypeInst ResType,
3258 unsigned OpAnyOrAll)
const {
3259 assert(
I.getNumOperands() == 3);
3260 assert(
I.getOperand(2).isReg());
3262 Register InputRegister =
I.getOperand(2).getReg();
3265 assert(InputType &&
"VReg has no type assigned");
3269 assert(ResVReg ==
I.getOperand(0).getReg());
3270 return BuildCOPY(ResVReg, InputRegister,
I);
3274 unsigned SpirvNotEqualId =
3275 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3277 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3282 IsBoolTy ? InputRegister
3290 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3292 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3309bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3310 SPIRVTypeInst ResType,
3311 MachineInstr &
I)
const {
3312 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3315bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3316 SPIRVTypeInst ResType,
3317 MachineInstr &
I)
const {
3318 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3322bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3323 SPIRVTypeInst ResType,
3324 MachineInstr &
I)
const {
3325 assert(
I.getNumOperands() == 4);
3326 assert(
I.getOperand(2).isReg());
3327 assert(
I.getOperand(3).isReg());
3329 [[maybe_unused]] SPIRVTypeInst VecType =
3334 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3335 "dot product requires either a vector of at least 2 components or"
3336 " the SPV_EXT_long vector extension.");
3338 [[maybe_unused]] SPIRVTypeInst EltType =
3347 .
addUse(
I.getOperand(2).getReg())
3348 .
addUse(
I.getOperand(3).getReg())
3353bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3354 SPIRVTypeInst ResType,
3357 assert(
I.getNumOperands() == 4);
3358 assert(
I.getOperand(2).isReg());
3359 assert(
I.getOperand(3).isReg());
3362 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3366 .
addUse(
I.getOperand(2).getReg())
3367 .
addUse(
I.getOperand(3).getReg())
3374bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3375 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3376 assert(
I.getNumOperands() == 4);
3377 assert(
I.getOperand(2).isReg());
3378 assert(
I.getOperand(3).isReg());
3382 Register Vec0 =
I.getOperand(2).getReg();
3383 Register Vec1 =
I.getOperand(3).getReg();
3387 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3396 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3397 "dot product requires either a vector of at least 2 components "
3398 "or the SPV_EXT_long_vector extension.");
3401 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3411 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3422 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3434bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3435 SPIRVTypeInst ResType,
3436 MachineInstr &
I)
const {
3438 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3441 .
addUse(
I.getOperand(2).getReg())
3446bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3447 SPIRVTypeInst ResType,
3448 MachineInstr &
I)
const {
3450 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3453 .
addUse(
I.getOperand(2).getReg())
3458bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3459 SPIRVTypeInst ResType,
3460 MachineInstr &
I)
const {
3463 Register Src =
I.getOperand(2).getReg();
3506bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3507 SPIRVTypeInst ResType,
3508 MachineInstr &
I)
const {
3510 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3513 .
addUse(
I.getOperand(2).getReg())
3518template <
bool Signed>
3519bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3520 SPIRVTypeInst ResType,
3521 MachineInstr &
I)
const {
3522 assert(
I.getNumOperands() == 5);
3523 assert(
I.getOperand(2).isReg());
3524 assert(
I.getOperand(3).isReg());
3525 assert(
I.getOperand(4).isReg());
3528 Register Acc =
I.getOperand(2).getReg();
3532 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3534 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3539 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3542 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3554template <
bool Signed>
3555bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3556 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3557 assert(
I.getNumOperands() == 5);
3558 assert(
I.getOperand(2).isReg());
3559 assert(
I.getOperand(3).isReg());
3560 assert(
I.getOperand(4).isReg());
3563 Register Acc =
I.getOperand(2).getReg();
3569 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3573 for (
unsigned i = 0; i < 4; i++) {
3596 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3616 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3631bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3632 SPIRVTypeInst ResType,
3633 MachineInstr &
I)
const {
3634 assert(
I.getNumOperands() == 3);
3635 assert(
I.getOperand(2).isReg());
3637 Register VZero = buildZerosValF(ResType,
I);
3638 Register VOne = buildOnesValF(ResType,
I);
3640 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3643 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3645 .
addUse(
I.getOperand(2).getReg())
3652bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3653 SPIRVTypeInst ResType,
3654 MachineInstr &
I)
const {
3655 assert(
I.getNumOperands() == 3);
3656 assert(
I.getOperand(2).isReg());
3658 Register InputRegister =
I.getOperand(2).getReg();
3660 auto &
DL =
I.getDebugLoc();
3663 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3670 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3672 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3680 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3685 if (NeedsConversion) {
3686 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3697bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3698 SPIRVTypeInst ResType,
3700 unsigned Opcode)
const {
3704 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3710 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3711 BMI.addUse(
I.getOperand(J).getReg());
3718bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3721 bool WithGroupSync)
const {
3723 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3725 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3727 assert(((Scope != SPIRV::Scope::Workgroup) ||
3728 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3729 "Workgroup Scope must set WorkGroupMemory semantic "
3730 "in Barrier instruction");
3732 assert(((Scope != SPIRV::Scope::Device) ||
3733 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3734 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3735 "Device Scope must set UniformMemory and ImageMemory semantic "
3736 "in Barrier instruction");
3742 if (WithGroupSync) {
3743 Register ExecReg = buildI32ConstantInEntryBlock(SPIRV::Scope::Workgroup,
I);
3747 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
3748 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
3750 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3754bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3755 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3760 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3761 SPIRV::OpGroupNonUniformBallot))
3766 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3771 .
addImm(SPIRV::GroupOperation::Reduce)
3778bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3779 SPIRVTypeInst ResType,
3780 MachineInstr &
I)
const {
3785 Register InputReg =
I.getOperand(2).getReg();
3790 bool IsVector = NumElems > 1 ||
3791 (InputType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
3805 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3806 SPIRV::OpGroupNonUniformAllEqual);
3811 ElementResults.
reserve(NumElems);
3813 for (
unsigned Idx = 0;
Idx < NumElems; ++
Idx) {
3826 ElemInput = Extracted;
3832 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3843 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3854bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3855 SPIRVTypeInst ResType,
3856 MachineInstr &
I)
const {
3858 assert(
I.getNumOperands() == 3);
3860 auto Op =
I.getOperand(2);
3870 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3872 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3873 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3894 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3898 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3905bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3906 SPIRVTypeInst ResType,
3908 bool IsUnsigned)
const {
3909 return selectWaveReduce(
3910 ResVReg, ResType,
I, IsUnsigned,
3911 [&](
Register InputRegister,
bool IsUnsigned) {
3912 const bool IsFloatTy =
3914 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3915 : SPIRV::OpGroupNonUniformSMax;
3916 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3920bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3921 SPIRVTypeInst ResType,
3923 bool IsUnsigned)
const {
3924 return selectWaveReduce(
3925 ResVReg, ResType,
I, IsUnsigned,
3926 [&](
Register InputRegister,
bool IsUnsigned) {
3927 const bool IsFloatTy =
3929 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3930 : SPIRV::OpGroupNonUniformSMin;
3931 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3935bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3936 SPIRVTypeInst ResType,
3937 MachineInstr &
I)
const {
3938 return selectWaveReduce(ResVReg, ResType,
I,
false,
3939 [&](
Register InputRegister,
bool IsUnsigned) {
3941 InputRegister, SPIRV::OpTypeFloat);
3942 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3943 : SPIRV::OpGroupNonUniformIAdd;
3947bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3948 SPIRVTypeInst ResType,
3949 MachineInstr &
I)
const {
3950 return selectWaveReduce(ResVReg, ResType,
I,
false,
3951 [&](
Register InputRegister,
bool IsUnsigned) {
3953 InputRegister, SPIRV::OpTypeFloat);
3954 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3955 : SPIRV::OpGroupNonUniformIMul;
3959template <
typename PickOpcodeFn>
3960bool SPIRVInstructionSelector::selectWaveReduce(
3961 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3962 PickOpcodeFn &&PickOpcode)
const {
3963 assert(
I.getNumOperands() == 3);
3964 assert(
I.getOperand(2).isReg());
3966 Register InputRegister =
I.getOperand(2).getReg();
3970 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3973 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3979 .
addImm(SPIRV::GroupOperation::Reduce)
3980 .
addUse(
I.getOperand(2).getReg())
3985bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3986 SPIRVTypeInst ResType,
3988 unsigned Opcode)
const {
3989 return selectWaveReduce(
3990 ResVReg, ResType,
I,
false,
3991 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3994bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3995 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3996 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3997 [&](
Register InputRegister,
bool IsUnsigned) {
3999 InputRegister, SPIRV::OpTypeFloat);
4001 ? SPIRV::OpGroupNonUniformFAdd
4002 : SPIRV::OpGroupNonUniformIAdd;
4006bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
4007 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4008 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
4009 [&](
Register InputRegister,
bool IsUnsigned) {
4011 InputRegister, SPIRV::OpTypeFloat);
4013 ? SPIRV::OpGroupNonUniformFMul
4014 : SPIRV::OpGroupNonUniformIMul;
4018template <
typename PickOpcodeFn>
4019bool SPIRVInstructionSelector::selectWaveExclusiveScan(
4020 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
4021 PickOpcodeFn &&PickOpcode)
const {
4022 assert(
I.getNumOperands() == 3);
4023 assert(
I.getOperand(2).isReg());
4025 Register InputRegister =
I.getOperand(2).getReg();
4029 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
4032 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
4038 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
4039 .
addUse(
I.getOperand(2).getReg())
4044bool SPIRVInstructionSelector::selectWaveMultiPrefixOp(
4045 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
unsigned Opcode,
4046 StringRef BuiltinName)
const {
4048 SPIRV::Extension::SPV_EXT_shader_subgroup_partitioned) &&
4050 SPIRV::Extension::SPV_NV_shader_subgroup_partitioned))
4051 return diagnoseUnsupported(
4052 I, Twine(BuiltinName) +
4053 " requires SPV_EXT_shader_subgroup_partitioned or "
4054 "SPV_NV_shader_subgroup_partitioned");
4056 assert(
I.getNumOperands() == 7);
4061 auto MaskMIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4062 TII.get(SPIRV::OpCompositeConstruct))
4065 for (
unsigned Op = 3;
Op != 7; ++
Op)
4066 MaskMIB.
addUse(
I.getOperand(
Op).getReg());
4069 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4074 .
addImm(SPIRV::GroupOperation::PartitionedExclusiveScanNV)
4075 .
addUse(
I.getOperand(2).getReg())
4081bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
4082 SPIRVTypeInst ResType,
4085 assert(
I.getNumOperands() == 3);
4086 assert(
I.getOperand(2).isReg());
4088 Register InputRegister =
I.getOperand(2).getReg();
4094 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
4105bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
4106 SPIRVTypeInst ResType,
4113 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
4118 : SPIRV::OpUConvert;
4120 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4123 ShiftOp = SPIRV::OpShiftRightLogicalV;
4128 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4129 TII.get(SPIRV::OpConstantComposite))
4132 for (
unsigned It = 0; It <
N; ++It)
4136 ShiftConst = CompositeReg;
4141 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
4146 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
4151 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
4156 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
4159bool SPIRVInstructionSelector::handle64BitOverflow(
4161 unsigned int Opcode,
4168 "handle64BitOverflow should only be used for integer types");
4170 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4172 MachineIRBuilder MIRBuilder(
I);
4174 SPIRVTypeInst I64x2Type =
4176 SPIRVTypeInst Vec2ResType =
4179 std::vector<Register> PartialRegs;
4181 unsigned CurrentComponent = 0;
4182 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4186 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4187 TII.get(SPIRV::OpVectorShuffle))
4192 .
addImm(CurrentComponent)
4193 .
addImm(CurrentComponent + 1);
4203 PartialRegs.push_back(SubVecReg);
4206 if (CurrentComponent != ComponentCount) {
4212 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4213 SPIRV::OpVectorExtractDynamic))
4222 PartialRegs.push_back(FinalElemResReg);
4226 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4227 SPIRV::OpCompositeConstruct);
4230bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4231 SPIRVTypeInst ResType,
4235 if (ComponentCount > 2)
4236 return handle64BitOverflow(
4237 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4239 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4241 MachineIRBuilder MIRBuilder(
I);
4245 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4249 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4254 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4261 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4262 TII.get(SPIRV::OpVectorShuffle))
4267 for (
unsigned J = 0; J < ComponentCount; ++J) {
4274 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4277bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4278 SPIRVTypeInst ResType,
4282 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4290bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4291 SPIRVTypeInst ResType,
4292 MachineInstr &
I)
const {
4293 Register OpReg =
I.getOperand(1).getReg();
4302 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4304 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4306 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4308 return SPIRVInstructionSelector::diagnoseUnsupported(
4309 I,
"G_BITREVERSE only support 16,32,64 bits.");
4313 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4324 unsigned AndOp = SPIRV::OpBitwiseAndS;
4325 unsigned OrOp = SPIRV::OpBitwiseOrS;
4326 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4327 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4328 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4330 AndOp = SPIRV::OpBitwiseAndV;
4331 OrOp = SPIRV::OpBitwiseOrV;
4332 ShlOp = SPIRV::OpShiftLeftLogicalV;
4333 ShrOp = SPIRV::OpShiftRightLogicalV;
4339 const unsigned Shift) ->
Register {
4342 (ResType->
getOpcode() != SPIRV::OpTypeVectorIdEXT ||
4349 Register MaskReg = CreateConst(Mask);
4350 Register ShiftReg = CreateConst(Shift);
4357 if (!selectOpWithSrcs(
T1, ResType,
I, {
Input, ShiftReg}, ShrOp) ||
4358 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4359 !selectOpWithSrcs(T3, ResType,
I, {
Input, MaskReg}, AndOp) ||
4360 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4361 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4370 while ((Shift >>= 1) > 0) {
4377 return BuildCOPY(ResVReg, Result,
I);
4380bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4381 SPIRVTypeInst ResType,
4382 MachineInstr &
I)
const {
4383 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4384 "G_FREEZE must define and use a register");
4385 Register OpReg =
I.getOperand(1).getReg();
4389 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4403 if (
Def->getOpcode() == TargetOpcode::COPY)
4406 switch (
Def->getOpcode()) {
4407 case SPIRV::ASSIGN_TYPE:
4408 if (MachineInstr *AssignToDef =
4410 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4411 Reg =
Def->getOperand(2).getReg();
4414 case SPIRV::OpUndef:
4415 Reg =
Def->getOperand(1).getReg();
4418 unsigned DestOpCode;
4420 DestOpCode = SPIRV::OpConstantNull;
4421 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4422 "static undef/poison lowered to OpConstantNull\n");
4424 DestOpCode = TargetOpcode::COPY;
4426 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4427 "skipped, lowered as a copy of the operand\n");
4429 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4430 .
addDef(
I.getOperand(0).getReg())
4438bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4439 SPIRVTypeInst ResType,
4440 MachineInstr &
I)
const {
4444 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4448 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4453 for (
unsigned i =
I.getNumExplicitDefs();
4454 i <
I.getNumExplicitOperands() && IsConst; ++i)
4459 return diagnoseUnsupported(
4460 I,
"There must be at least two constituent operands in a vector");
4465 for (
unsigned i =
I.getNumExplicitDefs();
4466 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4472 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4479 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4480 TII.get(IsConst ? SPIRV::OpConstantComposite
4481 : SPIRV::OpCompositeConstruct))
4484 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4485 MIB.
addUse(
I.getOperand(i).getReg());
4490bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4491 SPIRVTypeInst ResType,
4492 MachineInstr &
I)
const {
4496 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4501 unsigned OpIdx =
I.getNumExplicitDefs();
4502 if (!
I.getOperand(OpIdx).isReg())
4506 Register OpReg =
I.getOperand(OpIdx).getReg();
4510 return diagnoseUnsupported(
4511 I,
"There must be at least two constituent operands in a vector");
4514 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4515 TII.get(IsConst ? SPIRV::OpConstantComposite
4516 : SPIRV::OpCompositeConstruct))
4519 for (
unsigned i = 0; i <
N; ++i)
4525bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4526 SPIRVTypeInst ResType,
4527 MachineInstr &
I)
const {
4533 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4535 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4536 TII.get(SPIRV::OpCompositeConstruct))
4539 for (
unsigned OpIdx =
I.getNumExplicitDefs();
4540 OpIdx <
I.getNumExplicitOperands(); ++OpIdx)
4541 MIB.
addUse(
I.getOperand(OpIdx).getReg());
4546bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4547 SPIRVTypeInst ResType,
4548 MachineInstr &
I)
const {
4554 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4556 Opcode = SPIRV::OpDemoteToHelperInvocation;
4558 Opcode = SPIRV::OpKill;
4563 ToErase.eraseFromParent();
4572bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4573 SPIRVTypeInst ResType,
unsigned CmpOpc,
4574 MachineInstr &
I)
const {
4575 Register Cmp0 =
I.getOperand(2).getReg();
4576 Register Cmp1 =
I.getOperand(3).getReg();
4579 "CMP operands should have the same type");
4580 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4590bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4591 SPIRVTypeInst ResType,
4592 MachineInstr &
I)
const {
4596 Register CmpOperand =
I.getOperand(2).getReg();
4598 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4603 Register Op1 =
I.getOperand(3).getReg();
4607 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4612 I.getOperand(3).setReg(NewOp1);
4618 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4623Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4624 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4626 SPIRVTypeInst SpvI32Ty =
4628 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4633 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4634 MachineInstr *
MI =
nullptr;
4638 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4642 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4643 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4649 GR.
add(ConstInt,
MI);
4654bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4655 SPIRVTypeInst ResType,
4656 MachineInstr &
I)
const {
4665 for (
unsigned Idx = 0;
Idx < 2; ++
Idx) {
4667 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
4670 .
addUse(
I.getOperand(2 + Idx).getReg())
4677 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLogicalOr))
4684 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLogicalNot))
4692 return selectCmp(ResVReg, ResType, CmpOp,
I);
4695bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4696 SPIRVTypeInst ResType,
4697 MachineInstr &
I)
const {
4699 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4709 MachineIRBuilder MIRBuilder(
I);
4716 APFloat ConstVal(3.3219280948873623);
4720 APFloat::rmNearestTiesToEven, &LosesInfo);
4725 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
4727 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4728 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4730 if (!selectExtInst(ResVReg, ResType,
I,
4731 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4741Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4742 MachineInstr &
I)
const {
4750bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4756 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4764 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4767 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4768 Def->getOpcode() == SPIRV::OpConstantI)
4781 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4782 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4784 Intrinsic::spv_const_composite)) {
4785 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4786 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4787 if (!IsZero(
Def->getOperand(i).getReg()))
4796Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4797 MachineInstr &
I)
const {
4806Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4807 MachineInstr &
I)
const {
4817 SPIRVTypeInst ResType,
4818 MachineInstr &
I)
const {
4829 MachineInstr &
I)
const {
4834 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4835 TII.get(SPIRV::OpCompositeConstruct))
4838 for (
unsigned J = 0; J < NumElts; ++J)
4844bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4845 SPIRVTypeInst ResType,
4846 MachineInstr &
I)
const {
4847 Register SelectFirstArg =
I.getOperand(2).getReg();
4848 Register SelectSecondArg =
I.getOperand(3).getReg();
4857 Register CondReg =
I.getOperand(1).getReg();
4858 bool IsScalarBool = GR.
isScalarOfType(CondReg, SPIRV::OpTypeBool);
4866 CondReg = buildVectorSplat(CondReg, NumElts,
I);
4867 IsScalarBool =
false;
4870 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4871 }
else if (IsPtrTy) {
4872 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4874 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4877 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4878 "boolean condition");
4880 Opcode = SPIRV::OpSelectSFSCond;
4881 }
else if (IsPtrTy) {
4882 Opcode = SPIRV::OpSelectSPSCond;
4884 Opcode = SPIRV::OpSelectSISCond;
4887 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4899bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4900 SPIRVTypeInst ResType,
4902 MachineInstr &InsertAt,
4903 bool IsSigned)
const {
4905 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4906 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4907 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4909 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4921bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4922 SPIRVTypeInst ResType,
4923 MachineInstr &
I,
bool IsSigned,
4924 unsigned Opcode)
const {
4925 Register SrcReg =
I.getOperand(1).getReg();
4936 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4938 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4941bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4942 SPIRVTypeInst ResType, MachineInstr &
I,
4943 bool IsSigned)
const {
4944 Register SrcReg =
I.getOperand(1).getReg();
4946 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4950 if (ResType == SrcType)
4951 return BuildCOPY(ResVReg, SrcReg,
I);
4953 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4954 return selectUnOp(ResVReg, ResType,
I, Opcode);
4957bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4958 SPIRVTypeInst ResType,
4960 bool IsSigned)
const {
4961 MachineIRBuilder MIRBuilder(
I);
4962 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
4967 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4975 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4978 .
addUse(
I.getOperand(1).getReg())
4979 .
addUse(
I.getOperand(2).getReg())
4984 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4987 .
addUse(
I.getOperand(1).getReg())
4988 .
addUse(
I.getOperand(2).getReg())
4996 unsigned SelectOpcode =
4997 (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4999 ? SPIRV::OpSelectVIVCond
5000 : SPIRV::OpSelectSISCond;
5005 .
addUse(buildOnesVal(
true, ResType,
I))
5006 .
addUse(buildZerosVal(ResType,
I))
5013 .
addUse(buildOnesVal(
false, ResType,
I))
5018bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
5021 SPIRVTypeInst IntTy,
5022 SPIRVTypeInst BoolTy)
const {
5026 isVectorType(IntTy) ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
5028 Register One = buildOnesVal(
false, IntTy,
I);
5036 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
5045bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
5046 SPIRVTypeInst ResType,
5047 MachineInstr &
I)
const {
5048 Register IntReg =
I.getOperand(1).getReg();
5051 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
5052 if (ArgType == ResType)
5053 return BuildCOPY(ResVReg, IntReg,
I);
5055 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5056 return selectUnOp(ResVReg, ResType,
I, Opcode);
5059bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
5060 SPIRVTypeInst ResType,
5061 MachineInstr &
I)
const {
5062 unsigned Opcode =
I.getOpcode();
5063 unsigned TpOpcode = ResType->
getOpcode();
5065 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
5066 assert(Opcode == TargetOpcode::G_CONSTANT &&
5067 I.getOperand(1).getCImm()->isZero());
5068 MachineBasicBlock &DepMBB =
I.getMF()->front();
5071 }
else if (TpOpcode == SPIRV::OpTypeVectorIdEXT) {
5076 "Expected <1 x T> Vector!");
5077 if (Opcode == TargetOpcode::G_FCONSTANT)
5083 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
5091 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
5094bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
5095 SPIRVTypeInst ResType,
5096 MachineInstr &
I)
const {
5097 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5104bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
5105 SPIRVTypeInst ResType,
5106 MachineInstr &
I)
const {
5108 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
5112 .
addUse(
I.getOperand(3).getReg())
5114 .
addUse(
I.getOperand(2).getReg());
5115 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
5121bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
5122 SPIRVTypeInst ResType,
5123 MachineInstr &
I)
const {
5124 Type *MaybeResTy =
nullptr;
5129 "Expected aggregate type for extractv instruction");
5131 SPIRV::AccessQualifier::ReadWrite,
false);
5135 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5138 .
addUse(
I.getOperand(2).getReg());
5139 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
5145bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
5146 SPIRVTypeInst ResType,
5147 MachineInstr &
I)
const {
5149 return selectInsertVal(ResVReg, ResType,
I);
5151 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
5154 .
addUse(
I.getOperand(2).getReg())
5155 .
addUse(
I.getOperand(3).getReg())
5156 .
addUse(
I.getOperand(4).getReg())
5161bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
5162 SPIRVTypeInst ResType,
5163 MachineInstr &
I)
const {
5165 return selectExtractVal(ResVReg, ResType,
I);
5167 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
5170 .
addUse(
I.getOperand(2).getReg())
5171 .
addUse(
I.getOperand(3).getReg())
5176bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
5177 SPIRVTypeInst ResType,
5178 MachineInstr &
I)
const {
5179 const bool IsGEPInBounds =
I.getOperand(2).getImm();
5182 const bool UseUntypedPointers =
5183 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
5188 if (UseUntypedPointers) {
5190 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
5191 : SPIRV::OpUntypedAccessChainKHR;
5193 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
5194 : SPIRV::OpUntypedPtrAccessChainKHR;
5203 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
5205 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
5206 : SPIRV::OpPtrAccessChain;
5211 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5216 if (UseUntypedPointers) {
5231 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5232 Def->getOperand(1).isReg())
5234 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5235 if (
const auto *GVar =
5238 SPIRV::AccessQualifier::ReadWrite,
5242 return diagnoseUnsupported(
5243 I,
"could not deduce the base type of an untyped access chain");
5248 Res.addUse(BaseReg);
5250 const bool IsAccessChainOpcode =
5251 (Opcode == SPIRV::OpAccessChain ||
5252 Opcode == SPIRV::OpInBoundsAccessChain ||
5253 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5254 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5258 "Cannot translate GEP to OpAccessChain.");
5261 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5262 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5263 Res.addUse(
I.getOperand(i).getReg());
5264 Res.constrainAllUses(
TII,
TRI, RBI);
5273 if (Extract.
getOpcode() == SPIRV::OpCompositeExtract) {
5279 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5285 TII.get(SPIRV::OpCompositeInsert))
5299bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5301 unsigned Lim =
I.getNumExplicitOperands();
5302 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5303 Register OpReg =
I.getOperand(i).getReg();
5307 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5308 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5321 SPIRVTypeInst WrapType = OpType;
5322 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5324 SPIRV::StorageClass::CodeSectionINTEL) {
5326 SPIRV::StorageClass::Function,
I);
5333 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5334 TII.get(SPIRV::OpSpecConstantOp))
5337 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5339 GR.
add(OpDefine, MIB);
5345bool SPIRVInstructionSelector::selectDerivativeInst(
5346 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5347 const unsigned DPdOpCode)
const {
5350 if (!errorIfInstrOutsideShader(
I))
5356 Register SrcReg =
I.getOperand(2).getReg();
5361 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5364 .
addUse(
I.getOperand(2).getReg());
5366 MachineIRBuilder MIRBuilder(
I);
5369 if (componentCount != 1)
5377 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5382 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5387 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5395bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5396 SPIRVTypeInst ResType,
5397 MachineInstr &
I)
const {
5401 case Intrinsic::spv_load:
5402 return selectLoad(ResVReg, ResType,
I);
5403 case Intrinsic::spv_atomic_load:
5404 return selectAtomicLoad(ResVReg, ResType,
I);
5405 case Intrinsic::spv_store:
5406 return selectStore(
I);
5407 case Intrinsic::spv_atomic_store:
5408 return selectAtomicStore(
I);
5409 case Intrinsic::spv_extractv:
5410 return selectExtractVal(ResVReg, ResType,
I);
5411 case Intrinsic::spv_insertv:
5412 return selectInsertVal(ResVReg, ResType,
I);
5413 case Intrinsic::spv_extractelt:
5414 return selectExtractElt(ResVReg, ResType,
I);
5415 case Intrinsic::spv_insertelt:
5416 return selectInsertElt(ResVReg, ResType,
I);
5417 case Intrinsic::spv_gep:
5418 return selectGEP(ResVReg, ResType,
I);
5419 case Intrinsic::spv_bitcast: {
5420 Register OpReg =
I.getOperand(2).getReg();
5421 SPIRVTypeInst OpType =
5425 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5427 case Intrinsic::spv_unref_global:
5428 case Intrinsic::spv_init_global: {
5434 Register GVarVReg =
MI->getOperand(0).getReg();
5435 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5440 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5442 MI->eraseFromParent();
5446 case Intrinsic::spv_undef: {
5447 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5453 case Intrinsic::spv_poison:
5454 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5459 case Intrinsic::spv_freeze:
5460 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5463 .
addUse(
I.getOperand(2).getReg())
5466 case Intrinsic::spv_named_boolean_spec_constant: {
5467 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5468 : SPIRV::OpSpecConstantFalse;
5470 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5471 .
addDef(
I.getOperand(0).getReg())
5474 unsigned SpecId =
I.getOperand(2).getImm();
5476 SPIRV::Decoration::SpecId, {SpecId});
5480 case Intrinsic::spv_const_composite: {
5482 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5488 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5490 std::function<bool(
Register)> HasSpecConstOperand =
5500 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5501 J < Def->getNumExplicitOperands(); ++J) {
5502 if (
Def->getOperand(J).isReg() &&
5503 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5509 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5510 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5511 : SPIRV::OpConstantComposite;
5512 unsigned ContinuedOpc = HasSpecConst
5513 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5514 : SPIRV::OpConstantCompositeContinuedINTEL;
5515 MachineIRBuilder MIR(
I);
5517 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5519 for (
auto *Instr : Instructions) {
5520 Instr->setDebugLoc(
I.getDebugLoc());
5525 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5532 case Intrinsic::spv_assign_name: {
5533 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5534 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5535 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5536 i <
I.getNumExplicitOperands(); ++i) {
5537 MIB.
addImm(
I.getOperand(i).getImm());
5542 case Intrinsic::spv_switch: {
5543 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5544 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5545 if (
I.getOperand(i).isReg())
5546 MIB.
addReg(
I.getOperand(i).getReg());
5547 else if (
I.getOperand(i).isCImm())
5548 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5549 else if (
I.getOperand(i).isMBB())
5550 MIB.
addMBB(
I.getOperand(i).getMBB());
5557 case Intrinsic::spv_loop_merge: {
5558 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5559 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5560 if (
I.getOperand(i).isMBB())
5561 MIB.
addMBB(
I.getOperand(i).getMBB());
5568 case Intrinsic::spv_loop_control_intel: {
5570 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5571 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5576 case Intrinsic::spv_selection_merge: {
5578 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5579 assert(
I.getOperand(1).isMBB() &&
5580 "operand 1 to spv_selection_merge must be a basic block");
5581 MIB.
addMBB(
I.getOperand(1).getMBB());
5582 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5586 case Intrinsic::spv_cmpxchg:
5587 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5588 case Intrinsic::spv_unreachable:
5589 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5592 case Intrinsic::spv_abort:
5593 return selectAbort(
I);
5594 case Intrinsic::spv_alloca:
5595 return selectFrameIndex(ResVReg, ResType,
I);
5596 case Intrinsic::spv_alloca_array:
5597 return selectAllocaArray(ResVReg, ResType,
I);
5598 case Intrinsic::spv_assume:
5600 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5601 .
addUse(
I.getOperand(1).getReg())
5606 case Intrinsic::spv_expect:
5608 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5611 .
addUse(
I.getOperand(2).getReg())
5612 .
addUse(
I.getOperand(3).getReg())
5617 case Intrinsic::arithmetic_fence:
5618 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5619 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5622 .
addUse(
I.getOperand(2).getReg())
5626 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5628 case Intrinsic::spv_thread_id:
5634 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5636 case Intrinsic::spv_thread_id_in_group:
5642 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5644 case Intrinsic::spv_group_id:
5650 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5652 case Intrinsic::spv_flattened_thread_id_in_group:
5659 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5661 case Intrinsic::spv_workgroup_size:
5662 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5664 case Intrinsic::spv_global_size:
5665 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5667 case Intrinsic::spv_global_offset:
5668 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5670 case Intrinsic::spv_num_workgroups:
5671 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5673 case Intrinsic::spv_subgroup_size:
5674 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5676 case Intrinsic::spv_num_subgroups:
5677 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5679 case Intrinsic::spv_subgroup_id:
5680 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5681 case Intrinsic::spv_subgroup_local_invocation_id:
5682 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5683 ResVReg, ResType,
I);
5684 case Intrinsic::spv_subgroup_max_size:
5685 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5687 case Intrinsic::spv_fdot:
5688 return selectFloatDot(ResVReg, ResType,
I);
5689 case Intrinsic::spv_udot:
5690 case Intrinsic::spv_sdot:
5691 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5693 return selectIntegerDot(ResVReg, ResType,
I,
5694 IID == Intrinsic::spv_sdot);
5695 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5696 case Intrinsic::spv_dot4add_i8packed:
5697 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5699 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5700 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5701 case Intrinsic::spv_dot4add_u8packed:
5702 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5704 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5705 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5706 case Intrinsic::spv_all:
5707 return selectAll(ResVReg, ResType,
I);
5708 case Intrinsic::spv_any:
5709 return selectAny(ResVReg, ResType,
I);
5710 case Intrinsic::spv_distance:
5711 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5712 case Intrinsic::spv_lerp:
5713 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5714 case Intrinsic::spv_length:
5715 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5716 case Intrinsic::spv_degrees:
5717 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5718 case Intrinsic::spv_faceforward:
5719 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5720 case Intrinsic::spv_frac:
5721 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5722 case Intrinsic::spv_isinf:
5723 return selectOpIsInf(ResVReg, ResType,
I);
5724 case Intrinsic::spv_isnan:
5725 return selectOpIsNan(ResVReg, ResType,
I);
5726 case Intrinsic::spv_isfinite:
5727 return selectOpIsFinite(ResVReg, ResType,
I);
5728 case Intrinsic::spv_isnormal:
5729 return selectOpIsNormal(ResVReg, ResType,
I);
5730 case Intrinsic::spv_normalize:
5731 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5732 case Intrinsic::spv_refract:
5733 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5734 case Intrinsic::spv_reflect:
5735 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5736 case Intrinsic::spv_rsqrt:
5737 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5738 case Intrinsic::spv_sign:
5739 return selectSign(ResVReg, ResType,
I);
5740 case Intrinsic::spv_smoothstep:
5741 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5742 case Intrinsic::spv_firstbituhigh:
5743 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5744 case Intrinsic::spv_firstbitshigh:
5745 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5746 case Intrinsic::spv_firstbitlow:
5747 return selectFirstBitLow(ResVReg, ResType,
I);
5748 case Intrinsic::spv_all_memory_barrier:
5749 return selectBarrierInst(
I, SPIRV::Scope::Device,
5750 SPIRV::MemorySemantics::UniformMemory |
5751 SPIRV::MemorySemantics::ImageMemory |
5752 SPIRV::MemorySemantics::WorkgroupMemory,
5754 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5755 return selectBarrierInst(
I, SPIRV::Scope::Device,
5756 SPIRV::MemorySemantics::UniformMemory |
5757 SPIRV::MemorySemantics::ImageMemory |
5758 SPIRV::MemorySemantics::WorkgroupMemory,
5760 case Intrinsic::spv_device_memory_barrier:
5761 return selectBarrierInst(
I, SPIRV::Scope::Device,
5762 SPIRV::MemorySemantics::UniformMemory |
5763 SPIRV::MemorySemantics::ImageMemory,
5765 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5766 return selectBarrierInst(
I, SPIRV::Scope::Device,
5767 SPIRV::MemorySemantics::UniformMemory |
5768 SPIRV::MemorySemantics::ImageMemory,
5770 case Intrinsic::spv_group_memory_barrier:
5771 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5772 SPIRV::MemorySemantics::WorkgroupMemory,
5774 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5775 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5776 SPIRV::MemorySemantics::WorkgroupMemory,
5778 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5779 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5780 SPIRV::StorageClass::StorageClass ResSC =
5783 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5784 "from the Generic storage class");
5785 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5793 case Intrinsic::spv_lifetime_start:
5794 case Intrinsic::spv_lifetime_end: {
5795 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5796 : SPIRV::OpLifetimeStop;
5797 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5798 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5807 case Intrinsic::spv_saturate:
5808 return selectSaturate(ResVReg, ResType,
I);
5809 case Intrinsic::spv_nclamp:
5810 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5811 case Intrinsic::spv_uclamp:
5812 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5813 case Intrinsic::spv_sclamp:
5814 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5815 case Intrinsic::spv_subgroup_prefix_bit_count:
5816 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5817 case Intrinsic::spv_wave_active_countbits:
5818 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5819 case Intrinsic::spv_wave_all_equal:
5820 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5821 case Intrinsic::spv_wave_all:
5822 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5823 case Intrinsic::spv_wave_any:
5824 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5825 case Intrinsic::spv_subgroup_ballot:
5826 return selectWaveOpInst(ResVReg, ResType,
I,
5827 SPIRV::OpGroupNonUniformBallot);
5828 case Intrinsic::spv_wave_is_first_lane:
5829 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5830 case Intrinsic::spv_wave_reduce_or:
5831 return selectWaveReduceOp(ResVReg, ResType,
I,
5832 SPIRV::OpGroupNonUniformBitwiseOr);
5833 case Intrinsic::spv_wave_reduce_xor:
5834 return selectWaveReduceOp(ResVReg, ResType,
I,
5835 SPIRV::OpGroupNonUniformBitwiseXor);
5836 case Intrinsic::spv_wave_reduce_and:
5837 return selectWaveReduceOp(ResVReg, ResType,
I,
5838 SPIRV::OpGroupNonUniformBitwiseAnd);
5839 case Intrinsic::spv_wave_reduce_umax:
5840 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5841 case Intrinsic::spv_wave_reduce_max:
5842 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5843 case Intrinsic::spv_wave_reduce_umin:
5844 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5845 case Intrinsic::spv_wave_reduce_min:
5846 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5847 case Intrinsic::spv_wave_reduce_sum:
5848 return selectWaveReduceSum(ResVReg, ResType,
I);
5849 case Intrinsic::spv_wave_product:
5850 return selectWaveReduceProduct(ResVReg, ResType,
I);
5851 case Intrinsic::spv_wave_readlane:
5852 return selectWaveOpInst(ResVReg, ResType,
I,
5853 SPIRV::OpGroupNonUniformShuffle);
5854 case Intrinsic::spv_wave_readlane_first:
5855 return selectWaveOpInst(ResVReg, ResType,
I,
5856 SPIRV::OpGroupNonUniformBroadcastFirst);
5857 case Intrinsic::spv_wave_prefix_sum:
5858 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5859 case Intrinsic::spv_subgroup_multi_prefix_bit_and:
5860 return selectWaveMultiPrefixOp(ResVReg, ResType,
I,
5861 SPIRV::OpGroupNonUniformBitwiseAnd,
5862 "WaveMultiPrefixBitAnd");
5863 case Intrinsic::spv_subgroup_multi_prefix_bit_or:
5864 return selectWaveMultiPrefixOp(ResVReg, ResType,
I,
5865 SPIRV::OpGroupNonUniformBitwiseOr,
5866 "WaveMultiPrefixBitOr");
5867 case Intrinsic::spv_subgroup_multi_prefix_bit_xor:
5868 return selectWaveMultiPrefixOp(ResVReg, ResType,
I,
5869 SPIRV::OpGroupNonUniformBitwiseXor,
5870 "WaveMultiPrefixBitXor");
5871 case Intrinsic::spv_subgroup_multi_prefix_product: {
5874 return selectWaveMultiPrefixOp(ResVReg, ResType,
I,
5875 IsFloatTy ? SPIRV::OpGroupNonUniformFMul
5876 : SPIRV::OpGroupNonUniformIMul,
5877 "WaveMultiPrefixProduct");
5879 case Intrinsic::spv_subgroup_multi_prefix_sum: {
5882 return selectWaveMultiPrefixOp(ResVReg, ResType,
I,
5883 IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
5884 : SPIRV::OpGroupNonUniformIAdd,
5885 "WaveMultiPrefixSum");
5887 case Intrinsic::spv_wave_prefix_product:
5888 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5889 case Intrinsic::spv_quad_read_across_x: {
5890 return selectQuadSwap(ResVReg, ResType,
I, 0);
5892 case Intrinsic::spv_quad_read_across_y: {
5893 return selectQuadSwap(ResVReg, ResType,
I, 1);
5895 case Intrinsic::spv_quad_read_across_diagonal: {
5896 return selectQuadSwap(ResVReg, ResType,
I, 2);
5898 case Intrinsic::spv_radians:
5899 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5903 case Intrinsic::instrprof_increment:
5904 case Intrinsic::instrprof_increment_step:
5905 case Intrinsic::instrprof_value_profile:
5908 case Intrinsic::spv_value_md:
5910 case Intrinsic::spv_resource_handlefrombinding: {
5911 return selectHandleFromBinding(ResVReg, ResType,
I);
5913 case Intrinsic::spv_resource_counterhandlefrombinding:
5914 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5915 case Intrinsic::spv_resource_updatecounter:
5916 return selectUpdateCounter(ResVReg, ResType,
I);
5917 case Intrinsic::spv_resource_store_typedbuffer: {
5918 return selectImageWriteIntrinsic(
I);
5920 case Intrinsic::spv_resource_load_typedbuffer: {
5921 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5923 case Intrinsic::spv_resource_load_level: {
5924 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5926 case Intrinsic::spv_resource_getdimensions_x:
5927 case Intrinsic::spv_resource_getdimensions_xy:
5928 case Intrinsic::spv_resource_getdimensions_xyz: {
5929 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5931 case Intrinsic::spv_resource_getdimensions_levels_x:
5932 case Intrinsic::spv_resource_getdimensions_levels_xy:
5933 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5934 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5936 case Intrinsic::spv_resource_getdimensions_ms_xy:
5937 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5938 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5940 case Intrinsic::spv_resource_calculate_lod:
5941 case Intrinsic::spv_resource_calculate_lod_unclamped:
5942 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5943 case Intrinsic::spv_resource_sample:
5944 case Intrinsic::spv_resource_sample_clamp:
5945 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5946 case Intrinsic::spv_resource_samplebias:
5947 case Intrinsic::spv_resource_samplebias_clamp:
5948 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5949 case Intrinsic::spv_resource_samplegrad:
5950 case Intrinsic::spv_resource_samplegrad_clamp:
5951 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5952 case Intrinsic::spv_resource_samplelevel:
5953 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5954 case Intrinsic::spv_resource_samplecmp:
5955 case Intrinsic::spv_resource_samplecmp_clamp:
5956 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5957 case Intrinsic::spv_resource_samplecmplevelzero:
5958 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5959 case Intrinsic::spv_resource_gather:
5960 case Intrinsic::spv_resource_gather_cmp:
5961 return selectGatherIntrinsic(ResVReg, ResType,
I);
5962 case Intrinsic::spv_resource_getbasepointer:
5963 case Intrinsic::spv_resource_getpointer: {
5964 return selectResourceGetPointer(ResVReg, ResType,
I);
5966 case Intrinsic::spv_pushconstant_getpointer: {
5967 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5969 case Intrinsic::spv_discard: {
5970 return selectDiscard(ResVReg, ResType,
I);
5972 case Intrinsic::spv_resource_nonuniformindex: {
5973 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5975 case Intrinsic::spv_unpackhalf2x16: {
5976 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5978 case Intrinsic::spv_packhalf2x16: {
5979 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5981 case Intrinsic::spv_ddx:
5982 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5983 case Intrinsic::spv_ddy:
5984 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5985 case Intrinsic::spv_ddx_coarse:
5986 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5987 case Intrinsic::spv_ddy_coarse:
5988 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5989 case Intrinsic::spv_ddx_fine:
5990 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5991 case Intrinsic::spv_ddy_fine:
5992 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5993 case Intrinsic::spv_fwidth:
5994 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5995 case Intrinsic::spv_masked_gather:
5996 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5997 return selectMaskedGather(ResVReg, ResType,
I);
5998 return diagnoseUnsupported(
5999 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
6000 case Intrinsic::spv_masked_scatter:
6001 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
6002 return selectMaskedScatter(
I);
6003 return diagnoseUnsupported(
6004 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
6005 case Intrinsic::returnaddress:
6006 case Intrinsic::frameaddress: {
6008 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
6015 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
6020bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
6021 SPIRVTypeInst ResType,
6022 MachineInstr &
I)
const {
6027bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
6028 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6030 assert(Intr.getIntrinsicID() ==
6031 Intrinsic::spv_resource_counterhandlefrombinding);
6034 Register MainHandleReg = Intr.getOperand(2).getReg();
6036 assert(MainHandleDef->getIntrinsicID() ==
6037 Intrinsic::spv_resource_handlefrombinding);
6041 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(),
MRI);
6042 Register IndexReg = MainHandleDef->getOperand(5).getReg();
6043 std::string CounterName =
6048 MachineIRBuilder MIRBuilder(
I);
6050 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
6052 ArraySize, IndexReg, CounterName, MIRBuilder);
6054 return BuildCOPY(ResVReg, CounterVarReg,
I);
6057bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
6058 SPIRVTypeInst ResType,
6059 MachineInstr &
I)
const {
6061 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
6063 Register CounterHandleReg = Intr.getOperand(2).getReg();
6064 int64_t IncrVal = Intr.getOperand(3).getImm();
6071 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
6072 assert(CounterVarPointeeType &&
6073 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
6074 "Counter variable must be a struct");
6076 SPIRV::StorageClass::StorageBuffer &&
6077 "Counter variable must be in the storage buffer storage class");
6079 "Counter variable must have exactly 1 member in the struct");
6080 const SPIRVTypeInst MemberType =
6083 "Counter variable struct must have a single i32 member");
6087 MachineIRBuilder MIRBuilder(
I);
6089 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
6092 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
6098 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6101 .
addUse(CounterHandleReg)
6107 Register Scope = buildI32ConstantInEntryBlock(SPIRV::Scope::Device,
I);
6109 buildI32ConstantInEntryBlock(SPIRV::MemorySemantics::None,
I);
6112 buildI32ConstantInEntryBlock(
static_cast<uint32_t
>(IncrVal),
I);
6115 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
6124 return BuildCOPY(ResVReg, AtomicRes,
I);
6132 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
6140bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
6141 SPIRVTypeInst ResType,
6142 MachineInstr &
I)
const {
6150 Register ImageReg =
I.getOperand(2).getReg();
6158 Register IdxReg =
I.getOperand(3).getReg();
6160 MachineInstr &Pos =
I;
6162 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
6166bool SPIRVInstructionSelector::generateSampleImage(
6169 DebugLoc Loc, MachineInstr &Pos)
const {
6180 if (!loadHandleBeforePosition(NewSamplerReg,
6186 MachineIRBuilder MIRBuilder(Pos);
6199 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
6200 ImOps.Lod.has_value();
6201 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
6202 : SPIRV::OpImageSampleImplicitLod;
6204 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
6205 : SPIRV::OpImageSampleDrefImplicitLod;
6214 MIB.
addUse(*ImOps.Compare);
6216 uint32_t ImageOperands = 0;
6218 ImageOperands |= SPIRV::ImageOperand::Bias;
6220 ImageOperands |= SPIRV::ImageOperand::Lod;
6221 if (ImOps.GradX && ImOps.GradY)
6222 ImageOperands |= SPIRV::ImageOperand::Grad;
6223 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
6225 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6228 "Non-constant offsets are not supported in sample instructions.");
6233 ImageOperands |= SPIRV::ImageOperand::MinLod;
6235 if (ImageOperands != 0) {
6236 MIB.
addImm(ImageOperands);
6237 if (ImageOperands & SPIRV::ImageOperand::Bias)
6239 if (ImageOperands & SPIRV::ImageOperand::Lod)
6241 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6242 MIB.
addUse(*ImOps.GradX);
6243 MIB.
addUse(*ImOps.GradY);
6246 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6247 MIB.
addUse(*ImOps.Offset);
6248 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6249 MIB.
addUse(*ImOps.MinLod);
6256bool SPIRVInstructionSelector::selectImageQuerySize(
6258 std::optional<Register> LodReg)
const {
6261 "ImageReg is not an image type.");
6263 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6265 unsigned NumComponents = 0;
6267 case SPIRV::Dim::DIM_1D:
6268 case SPIRV::Dim::DIM_Buffer:
6269 NumComponents =
IsArray ? 2 : 1;
6271 case SPIRV::Dim::DIM_2D:
6272 case SPIRV::Dim::DIM_Cube:
6273 case SPIRV::Dim::DIM_Rect:
6274 NumComponents =
IsArray ? 3 : 2;
6276 case SPIRV::Dim::DIM_3D:
6280 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6290 if (!LodReg && IsSampled && !IsMultisampled &&
6291 Dim != SPIRV::Dim::DIM_Buffer && Dim != SPIRV::Dim::DIM_Rect)
6294 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6295 SPIRVTypeInst ResType =
6300 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6310bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6311 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6312 Register ImageReg =
I.getOperand(2).getReg();
6319 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6322bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6323 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6324 Register ImageReg =
I.getOperand(2).getReg();
6333 Register LodReg =
I.getOperand(3).getReg();
6336 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6338 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6345 TII.get(SPIRV::OpImageQueryLevels))
6352 TII.get(SPIRV::OpCompositeConstruct))
6362bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6363 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6364 Register ImageReg =
I.getOperand(2).getReg();
6375 "OpImageQuerySamples requires a multisampled image");
6377 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6385 TII.get(SPIRV::OpImageQuerySamples))
6392 TII.get(SPIRV::OpCompositeConstruct))
6402bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6403 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6404 Register ImageReg =
I.getOperand(2).getReg();
6405 Register SamplerReg =
I.getOperand(3).getReg();
6406 Register CoordinateReg =
I.getOperand(4).getReg();
6422 if (!loadHandleBeforePosition(
6427 MachineIRBuilder MIRBuilder(
I);
6433 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6443 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6450 unsigned ExtractedIndex =
6452 Intrinsic::spv_resource_calculate_lod_unclamped
6456 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6457 TII.get(SPIRV::OpCompositeExtract))
6467bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6468 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6469 Register ImageReg =
I.getOperand(2).getReg();
6470 Register SamplerReg =
I.getOperand(3).getReg();
6471 Register CoordinateReg =
I.getOperand(4).getReg();
6472 ImageOperands ImOps;
6473 if (
I.getNumOperands() > 5)
6474 ImOps.Offset =
I.getOperand(5).getReg();
6475 if (
I.getNumOperands() > 6)
6476 ImOps.MinLod =
I.getOperand(6).getReg();
6477 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6478 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6481bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6482 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6483 Register ImageReg =
I.getOperand(2).getReg();
6484 Register SamplerReg =
I.getOperand(3).getReg();
6485 Register CoordinateReg =
I.getOperand(4).getReg();
6486 ImageOperands ImOps;
6487 ImOps.Bias =
I.getOperand(5).getReg();
6488 if (
I.getNumOperands() > 6)
6489 ImOps.Offset =
I.getOperand(6).getReg();
6490 if (
I.getNumOperands() > 7)
6491 ImOps.MinLod =
I.getOperand(7).getReg();
6492 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6493 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6496bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6497 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6498 Register ImageReg =
I.getOperand(2).getReg();
6499 Register SamplerReg =
I.getOperand(3).getReg();
6500 Register CoordinateReg =
I.getOperand(4).getReg();
6501 ImageOperands ImOps;
6502 ImOps.GradX =
I.getOperand(5).getReg();
6503 ImOps.GradY =
I.getOperand(6).getReg();
6504 if (
I.getNumOperands() > 7)
6505 ImOps.Offset =
I.getOperand(7).getReg();
6506 if (
I.getNumOperands() > 8)
6507 ImOps.MinLod =
I.getOperand(8).getReg();
6508 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6509 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6512bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6513 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6514 Register ImageReg =
I.getOperand(2).getReg();
6515 Register SamplerReg =
I.getOperand(3).getReg();
6516 Register CoordinateReg =
I.getOperand(4).getReg();
6517 ImageOperands ImOps;
6518 ImOps.Lod =
I.getOperand(5).getReg();
6519 if (
I.getNumOperands() > 6)
6520 ImOps.Offset =
I.getOperand(6).getReg();
6521 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6522 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6525bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6526 SPIRVTypeInst ResType,
6527 MachineInstr &
I)
const {
6528 Register ImageReg =
I.getOperand(2).getReg();
6529 Register SamplerReg =
I.getOperand(3).getReg();
6530 Register CoordinateReg =
I.getOperand(4).getReg();
6531 ImageOperands ImOps;
6532 ImOps.Compare =
I.getOperand(5).getReg();
6533 if (
I.getNumOperands() > 6)
6534 ImOps.Offset =
I.getOperand(6).getReg();
6535 if (
I.getNumOperands() > 7)
6536 ImOps.MinLod =
I.getOperand(7).getReg();
6537 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6538 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6541bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6542 SPIRVTypeInst ResType,
6543 MachineInstr &
I)
const {
6544 Register ImageReg =
I.getOperand(2).getReg();
6545 Register CoordinateReg =
I.getOperand(3).getReg();
6546 Register LodReg =
I.getOperand(4).getReg();
6548 ImageOperands ImOps;
6550 if (
I.getNumOperands() > 5)
6551 ImOps.Offset =
I.getOperand(5).getReg();
6563 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6564 I.getDebugLoc(),
I, &ImOps);
6567bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6568 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6569 Register ImageReg =
I.getOperand(2).getReg();
6570 Register SamplerReg =
I.getOperand(3).getReg();
6571 Register CoordinateReg =
I.getOperand(4).getReg();
6572 ImageOperands ImOps;
6573 ImOps.Compare =
I.getOperand(5).getReg();
6574 if (
I.getNumOperands() > 6)
6575 ImOps.Offset =
I.getOperand(6).getReg();
6578 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6579 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6582bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6583 SPIRVTypeInst ResType,
6584 MachineInstr &
I)
const {
6585 Register ImageReg =
I.getOperand(2).getReg();
6586 Register SamplerReg =
I.getOperand(3).getReg();
6587 Register CoordinateReg =
I.getOperand(4).getReg();
6590 "ImageReg is not an image type.");
6595 ComponentOrCompareReg =
I.getOperand(5).getReg();
6596 OffsetReg =
I.getOperand(6).getReg();
6599 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6603 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6604 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6605 Dim != SPIRV::Dim::DIM_Rect) {
6607 "Gather operations are only supported for 2D, Cube, and Rect images.");
6614 if (!loadHandleBeforePosition(
6619 MachineIRBuilder MIRBuilder(
I);
6620 SPIRVTypeInst SampledImageType =
6625 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6633 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6635 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6637 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6642 .
addUse(ComponentOrCompareReg);
6644 uint32_t ImageOperands = 0;
6645 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6646 if (Dim == SPIRV::Dim::DIM_Cube) {
6648 "Gather operations with offset are not supported for Cube images.");
6652 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6654 ImageOperands |= SPIRV::ImageOperand::Offset;
6658 if (ImageOperands != 0) {
6659 MIB.
addImm(ImageOperands);
6661 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6669bool SPIRVInstructionSelector::generateImageReadOrFetch(
6672 const ImageOperands *ImOps)
const {
6675 "ImageReg is not an image type.");
6677 bool IsSignedInteger =
6682 bool IsFetch = (SampledOp.getImm() == 1);
6684 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6685 uint32_t ImageOperandsMask = 0;
6686 if (IsSignedInteger)
6687 ImageOperandsMask |= 0x1000;
6689 if (IsFetch && ImOps) {
6691 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6692 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6694 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6696 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6700 if (ImageOperandsMask != 0) {
6701 MIB.
addImm(ImageOperandsMask);
6702 if (IsFetch && ImOps) {
6705 if (ImOps->Offset &&
6706 (ImageOperandsMask &
6707 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6708 MIB.
addUse(*ImOps->Offset);
6717 SPIRVTypeInst SampledType =
6720 SPIRVTypeInst ReadType =
6721 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6722 bool ReadTypeMatchesResult = ReadType == ResType;
6724 Register ReadReg = ReadTypeMatchesResult
6730 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6736 BMI.constrainAllUses(
TII,
TRI, RBI);
6738 if (ReadTypeMatchesResult)
6751 if (ResultSize == 1) {
6760 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6763bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6764 SPIRVTypeInst ResType,
6765 MachineInstr &
I)
const {
6766 Register ResourcePtr =
I.getOperand(2).getReg();
6768 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6777 MachineIRBuilder MIRBuilder(
I);
6782 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6788 if (
I.getNumExplicitOperands() > 3) {
6789 Register IndexReg =
I.getOperand(3).getReg();
6796bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6797 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6802bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6803 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6804 Register ObjReg =
I.getOperand(2).getReg();
6805 if (!BuildCOPY(ResVReg, ObjReg,
I))
6815 decorateUsesAsNonUniform(ResVReg);
6819void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6822 {NonUniformReg,
nullptr}};
6823 llvm::SmallSet<Register, 8> Visited;
6824 while (WorkList.
size() > 0) {
6827 if (!Visited.
insert(CurrentReg).second)
6830 bool IsDecorated =
false;
6832 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6833 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6839 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6841 if (ResultReg == CurrentReg)
6849 MachineInstr &InsertPt =
6852 SPIRV::Decoration::NonUniformEXT, {});
6857bool SPIRVInstructionSelector::extractSubvector(
6859 MachineInstr &InsertionPoint)
const {
6861 [[maybe_unused]]
uint64_t InputSize =
6864 [[maybe_unused]]
bool IsLongVectorEXT =
6866 assert((InputSize > 1 || IsLongVectorEXT) &&
"The input must be a vector.");
6867 assert((ResultSize > 1 || IsLongVectorEXT) &&
"The result must be a vector.");
6868 assert(ResultSize < InputSize &&
6869 "Cannot extract more element than there are in the input.");
6876 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6885 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6887 TII.get(SPIRV::OpCompositeConstruct))
6891 for (
Register ComponentReg : ComponentRegisters)
6892 MIB.
addUse(ComponentReg);
6897bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6898 MachineInstr &
I)
const {
6905 Register ImageReg =
I.getOperand(1).getReg();
6913 Register CoordinateReg =
I.getOperand(2).getReg();
6914 Register DataReg =
I.getOperand(3).getReg();
6917 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6925Register SPIRVInstructionSelector::buildPointerToResource(
6926 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6927 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6928 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6930 if (ArraySize == 1) {
6931 SPIRVTypeInst PtrType =
6934 "SpirvResType did not have an explicit layout.");
6939 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6940 SPIRVTypeInst VarPointerType =
6943 VarPointerType, Set,
Binding, Name, MIRBuilder);
6945 SPIRVTypeInst ResPointerType =
6958bool SPIRVInstructionSelector::selectFirstBitSet16(
6959 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6960 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6962 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6966 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6969bool SPIRVInstructionSelector::selectFirstBitSet32(
6971 unsigned BitSetOpcode)
const {
6972 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6975 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6982bool SPIRVInstructionSelector::selectFirstBitSet64(
6984 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6998 if (ComponentCount > 2) {
6999 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
7001 unsigned Opcode) ->
bool {
7002 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
7006 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
7010 MachineIRBuilder MIRBuilder(
I);
7012 BaseType, 2 * ComponentCount, MIRBuilder,
false);
7016 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
7022 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
7032 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
7033 SPIRV::OpVectorExtractDynamic))
7035 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
7036 SPIRV::OpVectorExtractDynamic))
7040 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
7041 TII.get(SPIRV::OpVectorShuffle))
7049 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
7055 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
7056 TII.get(SPIRV::OpVectorShuffle))
7064 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
7084 SelectOp = SPIRV::OpSelectSISCond;
7085 AddOp = SPIRV::OpIAddS;
7093 SelectOp = SPIRV::OpSelectVIVCond;
7094 AddOp = SPIRV::OpIAddV;
7100 Register RegSecondaryOffset = Reg0;
7104 if (SwapPrimarySide) {
7105 PrimaryReg = LowReg;
7106 SecondaryReg = HighReg;
7107 RegPrimaryOffset = Reg0;
7108 RegSecondaryOffset = Reg32;
7113 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
7114 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
7119 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
7120 SPIRV::OpINotEqual))
7127 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
7128 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
7133 if (SwapPrimarySide) {
7135 if (!selectOpWithSrcs(RegAdd, ResType,
I,
7136 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
7147 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
7148 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
7153 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
7154 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
7157 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
7161bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
7162 SPIRVTypeInst ResType,
7164 bool IsSigned)
const {
7166 Register OpReg =
I.getOperand(2).getReg();
7169 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
7170 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
7174 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7176 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7178 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7181 return diagnoseUnsupported(
7183 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
7187bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
7188 SPIRVTypeInst ResType,
7189 MachineInstr &
I)
const {
7191 Register OpReg =
I.getOperand(2).getReg();
7196 unsigned ExtendOpcode = SPIRV::OpUConvert;
7197 unsigned BitSetOpcode = GL::FindILsb;
7201 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7203 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7205 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7208 return diagnoseUnsupported(
I,
7209 "spv_firstbitlow only supports 16,32,64 bits.");
7213bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
7214 SPIRVTypeInst ResType,
7215 MachineInstr &
I)
const {
7220 bool UseUntypedPointers =
7221 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7222 unsigned Opcode = UseUntypedPointers
7223 ? SPIRV::OpUntypedVariableLengthArrayINTEL
7224 : SPIRV::OpVariableLengthArrayINTEL;
7226 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
7232 if (UseUntypedPointers) {
7235 "untyped variable length array result must have a recorded element "
7240 MIB.
addUse(
I.getOperand(2).getReg());
7244 unsigned Alignment =
I.getOperand(3).getImm();
7258 while (!Worklist.
empty()) {
7260 switch (
T->getOpcode()) {
7261 case SPIRV::OpTypeInt:
7262 case SPIRV::OpTypeFloat:
7263 case SPIRV::OpTypePointer:
7265 case SPIRV::OpTypeVector:
7266 case SPIRV::OpTypeVectorIdEXT:
7267 case SPIRV::OpTypeMatrix:
7268 case SPIRV::OpTypeArray: {
7269 Register OperandReg =
T->getOperand(1).getReg();
7273 case SPIRV::OpTypeStruct:
7274 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7275 Register OperandReg =
T->getOperand(Idx).getReg();
7290 Register TypeReg = Ty->getOperand(0).getReg();
7291 if (!Visited.
insert(TypeReg).second)
7294 switch (Ty->getOpcode()) {
7295 case SPIRV::OpTypePointer:
7296 if (Ty->getOperand(1).getImm() == SPIRV::StorageClass::StorageBuffer)
7300 case SPIRV::OpTypeArray:
7301 case SPIRV::OpTypeRuntimeArray:
7304 case SPIRV::OpTypeStruct:
7305 for (
unsigned I = 1;
I < Ty->getNumOperands(); ++
I)
7321bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7322 assert(
I.getNumExplicitOperands() == 2);
7324 Register MsgReg =
I.getOperand(1).getReg();
7326 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7329 return diagnoseUnsupported(
7331 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7332 "scalar, pointer, vector, matrix, or aggregate of such types)");
7335 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7342bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7353 uint32_t MsgVal = ~0
u;
7354 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7355 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7358 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7361 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7368bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7369 SPIRVTypeInst ResType,
7370 MachineInstr &
I)
const {
7377 bool UseUntypedPointers =
7378 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7380 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7383 MachineIRBuilder MIRBuilder(
I);
7386 .
addImm(SPIRV::Extension::SPV_KHR_variable_pointers);
7388 .
addImm(SPIRV::Capability::VariablePointersStorageBuffer);
7391 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7394 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7398 if (UseUntypedPointers) {
7402 return diagnoseUnsupported(
7403 I,
"could not deduce the data type of an untyped variable");
7409 unsigned Alignment =
I.getOperand(2).getImm();
7416bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7421 const MachineInstr *PrevI =
I.getPrevNode();
7423 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7427 .
addMBB(
I.getOperand(0).getMBB())
7432 .
addMBB(
I.getOperand(0).getMBB())
7437bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7448 const MachineInstr *NextI =
I.getNextNode();
7450 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7456 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7458 .
addUse(
I.getOperand(0).getReg())
7459 .
addMBB(
I.getOperand(1).getMBB())
7465bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7466 MachineInstr &
I)
const {
7468 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7470 const unsigned NumOps =
I.getNumOperands();
7471 for (
unsigned i = 1; i <
NumOps; i += 2) {
7472 MIB.
addUse(
I.getOperand(i + 0).getReg());
7473 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7479bool SPIRVInstructionSelector::selectGlobalValue(
7480 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7482 MachineIRBuilder MIRBuilder(
I);
7483 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7486 std::string GlobalIdent;
7488 unsigned &
ID = UnnamedGlobalIDs[GV];
7490 ID = UnnamedGlobalIDs.
size();
7491 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7517 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7524 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
7529 MachineInstrBuilder MIB1 =
7530 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7533 MachineInstrBuilder MIB2 =
7535 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7539 GR.
add(ConstVal, MIB2);
7547 MachineInstrBuilder MIB3 =
7548 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7551 GR.
add(ConstVal, MIB3);
7557 assert(NewReg != ResVReg);
7558 return BuildCOPY(ResVReg, NewReg,
I);
7568 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7571 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7577 SPIRVTypeInst ResType =
7581 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7586 if (
GlobalVar->isExternallyInitialized() &&
7587 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7588 constexpr unsigned ReadWriteINTEL = 3u;
7591 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7597bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7598 SPIRVTypeInst ResType,
7599 MachineInstr &
I)
const {
7601 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7609 MachineIRBuilder MIRBuilder(
I);
7614 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7617 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7619 .
add(
I.getOperand(1))
7633 APFloat::rmNearestTiesToEven, &LosesInfo);
7638 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
7648bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7649 SPIRVTypeInst ResType,
7650 MachineInstr &
I)
const {
7653 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7659 Register ExpReg =
I.getOperand(2).getReg();
7661 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7662 SPIRV::OpConvertSToF))
7664 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7671bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7672 SPIRVTypeInst ResType,
7673 MachineInstr &
I)
const {
7689 MachineIRBuilder MIRBuilder(
I);
7690 SPIRVTypeInst FloatType =
7694 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7707 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7708 const bool IsUntyped =
7709 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7711 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7712 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7713 : SPIRV::OpVariable))
7716 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7724 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7727 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7730 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7734 Register IntegralPartReg =
I.getOperand(1).getReg();
7737 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7747 assert(
false &&
"GLSL::Modf is deprecated.");
7758bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7759 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7760 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7761 MachineIRBuilder MIRBuilder(
I);
7762 const SPIRVTypeInst Vec3Ty =
7765 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7777 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7781 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
7787 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7794 assert(
I.getOperand(2).isReg());
7795 const uint32_t ThreadId =
foldImm(
I.getOperand(2),
MRI);
7799 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7810bool SPIRVInstructionSelector::loadBuiltinInputID(
7811 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7812 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7813 MachineIRBuilder MIRBuilder(
I);
7815 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7830 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7834 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7843SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7844 MachineInstr &
I)
const {
7845 MachineIRBuilder MIRBuilder(
I);
7856bool SPIRVInstructionSelector::loadHandleBeforePosition(
7857 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7858 MachineInstr &Pos)
const {
7861 Intrinsic::spv_resource_handlefrombinding);
7869 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7870 MachineIRBuilder MIRBuilder(HandleDef);
7871 SPIRVTypeInst VarType = ResType;
7872 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7874 if (IsStructuredBuffer) {
7883 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7886 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7887 ArraySize, IndexReg, Name, MIRBuilder);
7891 uint32_t LoadOpcode =
7892 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7902bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7903 MachineInstr &
I)
const {
7905 return diagnoseUnsupported(
7906 I,
"this instruction is only supported in shaders.");
7911InstructionSelector *
7915 return new SPIRVInstructionSelector(TM,
Subtarget, RBI);
MachineInstrBuilder & UseMI
MachineInstrBuilder MachineInstrBuilder & DefMI
const AArch64Subtarget & Subtarget
#define GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
This file declares a class to represent arbitrary precision floating point values and provide a varie...
static bool selectUnmergeValues(MachineInstrBuilder &MIB, const ARMBaseInstrInfo &TII, MachineRegisterInfo &MRI, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
static uint8_t SwapBits(uint8_t Val)
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
DXIL Resource Implicit Binding
Declares convenience wrapper classes for interpreting MachineInstr instances as specific generic oper...
const HexagonInstrInfo * TII
static MaybeAlign getAlign(Value *Ptr)
LLVMTypeRef LLVMIntType(unsigned NumBits)
const size_t AbstractManglingParser< Derived, Alloc >::NumOps
Loop::LoopBounds::Direction Direction
Register const TargetRegisterInfo * TRI
Promote Memory to Register
uint64_t IntrinsicInst * II
static StringRef getName(Value *V)
static unsigned getFCmpOpcode(CmpInst::Predicate Pred, unsigned Size)
MachineRegisterInfo & MRI
static bool isConcreteSPIRVType(SPIRVTypeInst Ty, const SPIRVGlobalRegistry &GR)
static APFloat getOneFP(const Type *LLVMFloatTy)
static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC)
static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg)
static bool mayApplyGenericSelection(unsigned Opcode)
static APFloat getZeroFP(const Type *LLVMFloatTy)
std::vector< std::pair< SPIRV::InstructionSet::InstructionSet, uint32_t > > ExtInstList
static bool intrinsicHasSideEffects(Intrinsic::ID ID)
static unsigned getBoolCmpOpcode(unsigned PredNum)
static unsigned getICmpOpcode(unsigned PredNum)
static bool isOpcodeWithNoSideEffects(unsigned Opcode)
static bool containsStorageBufferPointer(SPIRVTypeInst Ty, const SPIRVGlobalRegistry &GR, SmallSet< Register, 8 > &Visited)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef)
static unsigned getPtrCmpOpcode(unsigned Pred)
static void addMemoryOperands(const MachineMemOperand *MemOp, MachineInstrBuilder &MIB, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry &GR, std::optional< Align > AlignOverride=std::nullopt)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
static Register convertPtrToInt(Register Reg, LLT ConvTy, SPIRVTypeInst SpvType, LegalizerHelper &Helper, MachineRegisterInfo &MRI, SPIRVGlobalRegistry *GR)
#define SPIRV_BACKEND_SERVICE_FUN_NAME
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
This file defines the SmallSet class.
static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")
static ManagedStatic< cl::opt< FnT >, OptCreatorT > CallbackFunction
static const fltSemantics & IEEEsingle()
static const fltSemantics & BFloat()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
const fltSemantics & getSemantics() const
static APFloat getOne(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative One.
static APFloat getZero(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative Zero.
static APInt getAllOnes(unsigned numBits)
Return an APInt of a specified width with all bits set.
static APInt getSignMask(unsigned BitWidth)
Get the SignMask for a specific bit width.
Represent a constant reference to an array (0 or more elements consecutively in memory),...
BlockFrequencyInfo pass uses BlockFrequencyInfoImpl implementation to estimate IR basic block frequen...
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
static LLVM_ABI std::optional< GFConstant > getConstant(Register Const, const MachineRegisterInfo &MRI)
static LLVM_ABI std::optional< GIConstant > getConstant(Register Const, const MachineRegisterInfo &MRI)
Represents a call to an intrinsic.
Intrinsic::ID getIntrinsicID() const
unsigned getAddressSpace() const
Module * getParent()
Get the module that this global value is contained inside of...
@ InternalLinkage
Rename collisions when linking (static functions).
static LLVM_ABI IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
constexpr bool isValid() const
constexpr uint16_t getNumElements() const
Returns the number of elements in a vector LLT.
constexpr bool isVector() const
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
int getNumber() const
MachineBasicBlocks are uniquely numbered at the function level, unless they're not in a MachineFuncti...
LLVM_ABI iterator getFirstNonPHI()
Returns a pointer to the first instruction in this block that is not a PHINode instruction.
instr_iterator instr_end()
const MachineFunction * getParent() const
Return the MachineFunction containing this basic block.
MachineInstrBundleIterator< MachineInstr > iterator
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
Helper class to build MachineInstr.
MachineBasicBlock::iterator getInsertPt()
Current insertion point for new instructions.
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
MachineFunction & getMF()
Getter for the function we currently build.
MachineRegisterInfo * getMRI()
Getter for MRI.
void constrainAllUses(const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI) const
const MachineInstrBuilder & addUse(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register use operand.
const MachineInstrBuilder & addReg(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a new virtual register operand.
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
const MachineInstrBuilder & add(const MachineOperand &MO) const
const MachineInstrBuilder & addMBB(MachineBasicBlock *MBB, unsigned TargetFlags=0) const
const MachineInstrBuilder & addDef(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register definition operand.
const MachineInstrBuilder & setMIFlags(unsigned Flags) const
MachineInstr * getInstr() const
If conversion operators fail, use this method to get the MachineInstr explicitly.
Representation of each machine instruction.
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
const MachineBasicBlock * getParent() const
unsigned getNumOperands() const
Retuns the total number of operands.
LLVM_ABI unsigned getNumExplicitOperands() const
Returns the number of non-implicit operands.
LLVM_ABI unsigned getNumExplicitDefs() const
Returns the number of non-implicit definitions.
LLVM_ABI void substituteRegister(Register FromReg, Register ToReg, unsigned SubIdx, const TargetRegisterInfo &RegInfo)
Replace all occurrences of FromReg with ToReg:SubIdx, properly composing subreg indices where necessa...
LLVM_ABI void emitGenericError(const Twine &ErrMsg) const
LLVM_ABI const MachineFunction * getMF() const
Return the function that contains the basic block that this instruction belongs to.
const DebugLoc & getDebugLoc() const
Returns the debug location id of this MachineInstr.
const MachineOperand & getOperand(unsigned i) const
A description of a memory reference used in the backend.
@ MOVolatile
The memory access is volatile.
@ MONonTemporal
The memory access is non-temporal.
LLVM_ABI Align getAlign() const
Return the minimum known alignment in bytes of the actual memory reference.
bool isReg() const
isReg - Tests if this is a MO_Register operand.
MachineBasicBlock * getMBB() const
Register getReg() const
getReg - Returns the register number.
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
defusechain_instr_iterator< true, false, false, true > use_instr_iterator
use_instr_iterator/use_instr_begin/use_instr_end - Walk all uses of the specified register,...
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
LLVM_ABI LLVM_READONLY MachineInstr * getVRegDef(Register Reg) const
getVRegDef - Return the machine instr that defines the specified virtual register or null if none is ...
use_instr_iterator use_instr_begin(Register RegNo) const
bool use_nodbg_empty(Register RegNo) const
use_nodbg_empty - Return true if there are no non-Debug instructions using the specified register.
static def_instr_iterator def_instr_end()
defusechain_instr_iterator< false, true, false, true > def_instr_iterator
def_instr_iterator/def_instr_begin/def_instr_end - Walk all defs of the specified register,...
LLVM_ABI Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
def_instr_iterator def_instr_begin(Register RegNo) const
LLT getType(Register Reg) const
Get the low-level type of Reg or LLT{} if Reg is not a generic (target independent) virtual register.
bool hasOneUse(Register RegNo) const
hasOneUse - Return true if there is exactly one instruction using the specified register.
static use_instr_iterator use_instr_end()
iterator_range< use_instr_nodbg_iterator > use_nodbg_instructions(Register Reg) const
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
const MachineFunction & getMF() const
LLVM_ABI void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
LLVM_ABI Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
const TargetRegisterClass * getRegClassOrNull(Register Reg) const
Return the register class of Reg, or null if Reg has not been assigned a register class yet.
iterator_range< use_instr_iterator > use_instructions(Register Reg) const
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
LLVM_ABI void replaceRegWith(Register FromReg, Register ToReg)
replaceRegWith - Replace all instances of FromReg with ToReg in the machine function.
Analysis providing profile information.
Holds all the information related to register banks.
Wrapper class representing virtual and physical registers.
constexpr bool isValid() const
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
const SDValue & getOperand(unsigned i) const
bool isScalarOrVectorSigned(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateOpTypeSampledImage(SPIRVTypeInst ImageType, MachineIRBuilder &MIRBuilder)
void assignSPIRVTypeToVReg(SPIRVTypeInst Type, Register VReg, const MachineFunction &MF)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC, bool ForceTyped=false)
const TargetRegisterClass * getRegClass(SPIRVTypeInst SpvType) const
MachineInstr * getOrAddMemAliasingINTELInst(MachineIRBuilder &MIRBuilder, const MDNode *AliasingListMD)
bool isAggregateType(SPIRVTypeInst Type) const
unsigned getScalarOrVectorBitWidth(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst getOrCreateSPIRVVectorType(SPIRVTypeInst BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
Register buildGlobalVariable(Register Reg, SPIRVTypeInst BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
SPIRVTypeInst getResultType(Register VReg, MachineFunction *MF=nullptr)
unsigned getScalarOrVectorComponentCount(Register VReg) const
const Type * getTypeForSPIRVType(SPIRVTypeInst Ty) const
bool isBitcastCompatible(SPIRVTypeInst Type1, SPIRVTypeInst Type2) const
unsigned getPointerSize() const
Register getOrCreateConstFP(APFloat Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
LLT getRegType(SPIRVTypeInst SpvType) const
SPIRVTypeInst getOpTypeVoid(MachineIRBuilder &MIRBuilder)
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
Register getSPIRVTypeID(SPIRVTypeInst SpirvType) const
Register getOrCreateConstInt(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
bool findValueAttrs(const MachineInstr *Key, Type *&Ty, StringRef &Name)
SPIRVTypeInst retrieveScalarOrVectorIntType(SPIRVTypeInst Type) const
Register getOrCreateGlobalVariableWithBinding(SPIRVTypeInst VarType, uint32_t Set, uint32_t Binding, StringRef Name, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst changePointerStorageClass(SPIRVTypeInst PtrType, SPIRV::StorageClass::StorageClass SC, MachineInstr &I)
Register getOrCreateConstVector(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType=nullptr)
void addGlobalObject(const Value *V, const MachineFunction *MF, Register R)
SPIRVTypeInst getScalarOrVectorComponentType(SPIRVTypeInst Type) const
void recordFunctionPointer(const MachineOperand *MO, const Function *F)
SPIRVTypeInst getOrCreateSPIRVFloatType(unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII)
SPIRVTypeInst getPointeeType(SPIRVTypeInst PtrType)
SPIRVTypeInst getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
MachineFunction * setCurrentFunc(MachineFunction &MF)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType)
SPIRVTypeInst getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
Type * getDeducedGlobalValueType(const GlobalValue *Global)
Register getOrCreateUndef(MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
SPIRVTypeInst getUntypedPtrElementType(Register Reg) const
bool erase(const MachineInstr *MI)
bool add(SPIRV::IRHandle Handle, const MachineInstr *MI)
Register find(SPIRV::IRHandle Handle, const MachineFunction *MF)
bool isPhysicalSPIRV() const
bool isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const
bool canUseExtInstSet(SPIRV::InstructionSet::InstructionSet E) const
bool isLogicalSPIRV() const
bool canUseExtension(SPIRV::Extension::Extension E) const
bool isTypeIntOrFloat() const
bool isAnyTypeFloat() const
bool erase(PtrType Ptr)
Remove pointer from the set.
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
bool contains(ConstPtrType Ptr) const
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
SmallSet - This maintains a set of unique values, optimizing for the case when the set is small (less...
std::pair< const_iterator, bool > insert(const T &V)
insert - Insert an element into the set if it isn't already there.
reference emplace_back(ArgTypes &&... Args)
void reserve(size_type N)
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
Represent a constant reference to a string, i.e.
constexpr size_t size() const
Get the string size.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
@ HalfTyID
16-bit floating point type
@ FloatTyID
32-bit floating point type
@ BFloatTyID
16-bit floating point type (7-bit significand)
@ DoubleTyID
64-bit floating point type
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
bool isStructTy() const
True if this is an instance of StructType.
bool isAggregateType() const
Return true if the type is an aggregate type.
TypeID getTypeID() const
Return the type id for the type.
Value * getOperand(unsigned i) const
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
An efficient, type-erasing, non-owning reference to a callable.
self_iterator getIterator()
NodeTy * getNextNode()
Get the next node, or nullptr for the list tail.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
constexpr char Align[]
Key for Kernel::Arg::Metadata::mAlign.
constexpr char IsConst[]
Key for Kernel::Arg::Metadata::mIsConst.
constexpr std::underlying_type_t< E > Mask()
Get a bitmask with 1s in all places up to the high-order bit of E's largest value.
ElementType
The element type of an SRV or UAV resource.
constexpr uint64_t PointerSize
aarch64 pointer size.
Scope
Defines the scope in which this symbol should be visible: Default – Visible in the public interface o...
NodeAddr< DefNode * > Def
NodeAddr< InstrNode * > Instr
NodeAddr< UseNode * > Use
NodeAddr< FuncNode * > Func
BaseReg
Stack frame base register. Bit 0 of FREInfo.Info.
unsigned getOpcode(const VPValue *V)
Return the instruction opcode for the recipe defining V or 0 for unsupported recipes and VPValues not...
This is an optimization pass for GlobalISel generic memory operations.
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
void addStringImm(StringRef Str, MCInst &Inst)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineFunction &MF)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
uint32_t getMemSemanticsWithStorageClass(const Triple &TT, uint32_t OrderSem, uint32_t StorageClassSem)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
MachineInstr * getDef(const MachineOperand &MO, const MachineRegisterInfo *MRI)
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
LLVM_ABI void salvageDebugInfo(const MachineRegisterInfo &MRI, MachineInstr &MI)
Assuming the instruction MI is going to be deleted, attempt to salvage debug users of MI by writing t...
LLVM_ABI void constrainSelectedInstRegOperands(MachineInstr &I, const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
Mutate the newly-selected instruction I to constrain its (possibly generic) virtual register operands...
bool isPreISelGenericOpcode(unsigned Opcode)
Check whether the given Opcode is a generic opcode that is not supposed to appear after ISel.
iterator_range< T > make_range(T x, T y)
Convenience function for iterating over sub-ranges.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, ArrayRef< uint32_t > DecArgs, StringRef StrImm)
LLVM_ABI bool isNullOrNullSplat(const MachineInstr &MI, const MachineRegisterInfo &MRI, bool AllowUndefs=false)
Return true if the value is a constant 0 integer or a splatted vector of a constant 0 integer (with n...
SPIRV::Scope::Scope getMemScope(const Triple &TT, LLVMContext &Ctx, SyncScope::ID Id)
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
bool isVectorType(SPIRVTypeInst SPVTy)
RelativeUniformCounterPtr ValuesPtrExpr VTableAddr Value
void buildOpName(Register Target, StringRef Name, MachineIRBuilder &MIRBuilder)
bool any_of(R &&range, UnaryPredicate P)
Provide wrappers to std::any_of which take ranges instead of having to pass begin/end explicitly.
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
Type * toTypedPointer(Type *Ty)
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
constexpr bool isGenericCastablePtr(SPIRV::StorageClass::StorageClass SC)
class LLVM_GSL_OWNER SmallVector
Forward declaration of SmallVector so that calculateSmallVectorDefaultInlinedElements can reference s...
MachineInstr * passCopy(MachineInstr *Def, const MachineRegisterInfo *MRI)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
std::optional< SPIRV::LinkageType::LinkageType > getSpirvLinkageTypeFor(const SPIRVSubtarget &ST, const GlobalValue &GV)
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
SPIRV::StorageClass::StorageClass addressSpaceToStorageClass(unsigned AddrSpace, const SPIRVSubtarget &STI)
AtomicOrdering
Atomic ordering for LLVM's memory model.
InstructionSelector * createSPIRVInstructionSelector(const SPIRVTargetMachine &TM, const SPIRVSubtarget &Subtarget, const RegisterBankInfo &RBI)
std::string getStringValueFromReg(Register Reg, MachineRegisterInfo &MRI)
int64_t foldImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
DWARFExpression::Operation Op
ArrayRef(const T &OneElt) -> ArrayRef< T >
MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)
constexpr unsigned BitWidth
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
bool hasInitializer(const GlobalVariable *GV)
bool isSpvIntrinsic(const MachineInstr &MI, Intrinsic::ID IntrinsicID)
MachineInstr * getVRegDef(MachineRegisterInfo &MRI, Register Reg)
SPIRV::MemorySemantics::MemorySemantics getMemSemantics(AtomicOrdering Ord)
std::string getLinkStringForBuiltIn(SPIRV::BuiltIn::BuiltIn BuiltInValue)
LLVM_ABI bool isTriviallyDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
Check whether an instruction MI is dead: it only defines dead virtual registers, and doesn't have oth...
MCRegisterClass TargetRegisterClass