34#include "llvm/IR/IntrinsicsSPIRV.h"
40#define DEBUG_TYPE "spirv-isel"
47 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
52 std::optional<Register> Bias;
53 std::optional<Register>
Offset;
54 std::optional<Register> MinLod;
55 std::optional<Register> GradX;
56 std::optional<Register> GradY;
57 std::optional<Register> Lod;
58 std::optional<Register> Compare;
65 bool IsScalar =
false;
68llvm::SPIRV::SelectionControl::SelectionControl
69getSelectionOperandForImm(
int Imm) {
71 return SPIRV::SelectionControl::Flatten;
73 return SPIRV::SelectionControl::DontFlatten;
75 return SPIRV::SelectionControl::None;
79#define GET_GLOBALISEL_PREDICATE_BITSET
80#include "SPIRVGenGlobalISel.inc"
81#undef GET_GLOBALISEL_PREDICATE_BITSET
108#define GET_GLOBALISEL_PREDICATES_DECL
109#include "SPIRVGenGlobalISel.inc"
110#undef GET_GLOBALISEL_PREDICATES_DECL
112#define GET_GLOBALISEL_TEMPORARIES_DECL
113#include "SPIRVGenGlobalISel.inc"
114#undef GET_GLOBALISEL_TEMPORARIES_DECL
138 unsigned BitSetOpcode)
const;
142 unsigned BitSetOpcode)
const;
146 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
153 unsigned Opcode)
const;
156 unsigned Opcode)
const;
178 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
189 unsigned OpType)
const;
252 unsigned Opcode)
const;
256 unsigned Opcode)
const;
260 unsigned Opcode)
const;
262 template <
bool Signed>
265 template <
bool Signed>
272 template <
typename PickOpcodeFn>
275 PickOpcodeFn &&PickOpcode)
const;
292 template <
typename PickOpcodeFn>
295 PickOpcodeFn &&PickOpcode)
const;
313 bool IsSigned)
const;
315 bool IsSigned,
unsigned Opcode)
const;
317 bool IsSigned)
const;
323 bool IsSigned)
const;
364 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
365 bool useMISrc =
true,
367 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
368 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
369 bool useMISrc =
true,
371 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
372 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
373 bool setMIFlags =
true,
bool useMISrc =
true,
375 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
376 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
377 bool useMISrc =
true,
380 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
381 MachineInstr &
I)
const;
383 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
384 MachineInstr &
I)
const;
386 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
387 MachineInstr &
I)
const;
389 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
390 MachineInstr &
I,
unsigned Opcode)
const;
392 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
393 bool WithGroupSync)
const;
395 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
396 MachineInstr &
I)
const;
398 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
399 MachineInstr &
I)
const;
403 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
404 MachineInstr &
I)
const;
406 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
407 MachineInstr &
I)
const;
409 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
411 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
412 MachineInstr &
I)
const;
413 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
414 SPIRVTypeInst ResType,
415 MachineInstr &
I)
const;
416 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
417 MachineInstr &
I)
const;
420 std::optional<Register> LodReg = std::nullopt)
const;
421 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
422 MachineInstr &
I)
const;
423 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
424 MachineInstr &
I)
const;
425 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
426 MachineInstr &
I)
const;
427 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
428 MachineInstr &
I)
const;
429 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
430 MachineInstr &
I)
const;
431 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
432 MachineInstr &
I)
const;
433 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
434 MachineInstr &
I)
const;
435 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
436 SPIRVTypeInst ResType,
437 MachineInstr &
I)
const;
438 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
439 MachineInstr &
I)
const;
440 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
441 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
442 MachineInstr &
I)
const;
443 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
444 MachineInstr &
I)
const;
445 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
446 MachineInstr &
I)
const;
447 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
450 MachineInstr &
I)
const;
451 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
452 MachineInstr &
I)
const;
453 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
454 MachineInstr &
I)
const;
455 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
458 MachineInstr &
I,
const unsigned DPdOpCode)
const;
460 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
461 SPIRVTypeInst ResType =
nullptr)
const;
462 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
463 SPIRVTypeInst ResType =
nullptr)
const;
465 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
466 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
467 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
469 MachineInstr &
I)
const;
470 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
472 bool wrapIntoSpecConstantOp(MachineInstr &
I,
475 Register getUcharPtrTypeReg(MachineInstr &
I,
476 SPIRV::StorageClass::StorageClass SC)
const;
477 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
479 uint32_t Opcode)
const;
480 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
481 SPIRVTypeInst SrcPtrTy)
const;
482 Register buildPointerToResource(SPIRVTypeInst ResType,
483 SPIRV::StorageClass::StorageClass SC,
484 uint32_t Set, uint32_t
Binding,
485 uint32_t ArraySize,
Register IndexReg,
487 MachineIRBuilder MIRBuilder)
const;
488 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
489 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
490 Register &ReadReg, MachineInstr &InsertionPoint)
const;
491 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
494 const ImageOperands *ImOps =
nullptr)
const;
495 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
497 Register CoordinateReg,
const ImageOperands &ImOps,
500 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
501 Register ResVReg, SPIRVTypeInst ResType,
502 MachineInstr &
I)
const;
503 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
504 Register ResVReg, SPIRVTypeInst ResType,
505 MachineInstr &
I)
const;
506 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
507 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
508 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
509 void errorIfInstrOutsideShader(MachineInstr &
I)
const;
511 std::optional<SplitParts> splitEvenOddLanes(
Register PopCountReg,
512 unsigned ComponentCount,
514 SPIRVTypeInst I32Type)
const;
517 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
518 Register SrcReg,
unsigned int Opcode,
519 std::function<
bool(
Register, SPIRVTypeInst,
520 MachineInstr &,
Register,
unsigned)>
524bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
526 if (
TET->getTargetExtName() ==
"spirv.Image") {
529 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
530 return TET->getTypeParameter(0)->isIntegerTy();
534#define GET_GLOBALISEL_IMPL
535#include "SPIRVGenGlobalISel.inc"
536#undef GET_GLOBALISEL_IMPL
542 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
545#include
"SPIRVGenGlobalISel.inc"
548#include
"SPIRVGenGlobalISel.inc"
560 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
564void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
565 if (HasVRegsReset == &MF)
580 for (
const auto &
MBB : MF) {
581 for (
const auto &
MI :
MBB) {
584 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
588 LLT DstType = MRI.
getType(DstReg);
590 LLT SrcType = MRI.
getType(SrcReg);
591 if (DstType != SrcType)
596 if (DstRC != SrcRC && SrcRC)
608 while (!Stack.empty()) {
613 switch (
MI->getOpcode()) {
614 case TargetOpcode::G_INTRINSIC:
615 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
616 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
619 if (IntrID != Intrinsic::spv_const_composite &&
620 IntrID != Intrinsic::spv_undef)
624 case TargetOpcode::G_BUILD_VECTOR:
625 case TargetOpcode::G_SPLAT_VECTOR:
627 i < OpDef->getNumOperands(); i++) {
632 Stack.push_back(OpNestedDef);
635 case TargetOpcode::G_CONSTANT:
636 case TargetOpcode::G_FCONSTANT:
637 case TargetOpcode::G_IMPLICIT_DEF:
638 case SPIRV::OpConstantTrue:
639 case SPIRV::OpConstantFalse:
640 case SPIRV::OpConstantI:
641 case SPIRV::OpConstantF:
642 case SPIRV::OpConstantComposite:
643 case SPIRV::OpConstantCompositeContinuedINTEL:
644 case SPIRV::OpConstantSampler:
645 case SPIRV::OpConstantNull:
647 case SPIRV::OpConstantFunctionPointerINTEL:
674 case Intrinsic::spv_all:
675 case Intrinsic::spv_alloca:
676 case Intrinsic::spv_any:
677 case Intrinsic::spv_bitcast:
678 case Intrinsic::spv_const_composite:
679 case Intrinsic::spv_cross:
680 case Intrinsic::spv_degrees:
681 case Intrinsic::spv_distance:
682 case Intrinsic::spv_extractelt:
683 case Intrinsic::spv_extractv:
684 case Intrinsic::spv_faceforward:
685 case Intrinsic::spv_fdot:
686 case Intrinsic::spv_firstbitlow:
687 case Intrinsic::spv_firstbitshigh:
688 case Intrinsic::spv_firstbituhigh:
689 case Intrinsic::spv_frac:
690 case Intrinsic::spv_gep:
691 case Intrinsic::spv_global_offset:
692 case Intrinsic::spv_global_size:
693 case Intrinsic::spv_group_id:
694 case Intrinsic::spv_insertelt:
695 case Intrinsic::spv_insertv:
696 case Intrinsic::spv_isinf:
697 case Intrinsic::spv_isnan:
698 case Intrinsic::spv_lerp:
699 case Intrinsic::spv_length:
700 case Intrinsic::spv_normalize:
701 case Intrinsic::spv_num_subgroups:
702 case Intrinsic::spv_num_workgroups:
703 case Intrinsic::spv_ptrcast:
704 case Intrinsic::spv_radians:
705 case Intrinsic::spv_reflect:
706 case Intrinsic::spv_refract:
707 case Intrinsic::spv_resource_getbasepointer:
708 case Intrinsic::spv_resource_getpointer:
709 case Intrinsic::spv_resource_handlefrombinding:
710 case Intrinsic::spv_resource_handlefromimplicitbinding:
711 case Intrinsic::spv_resource_nonuniformindex:
712 case Intrinsic::spv_resource_sample:
713 case Intrinsic::spv_rsqrt:
714 case Intrinsic::spv_saturate:
715 case Intrinsic::spv_sdot:
716 case Intrinsic::spv_sign:
717 case Intrinsic::spv_smoothstep:
718 case Intrinsic::spv_step:
719 case Intrinsic::spv_subgroup_id:
720 case Intrinsic::spv_subgroup_local_invocation_id:
721 case Intrinsic::spv_subgroup_max_size:
722 case Intrinsic::spv_subgroup_size:
723 case Intrinsic::spv_thread_id:
724 case Intrinsic::spv_thread_id_in_group:
725 case Intrinsic::spv_udot:
726 case Intrinsic::spv_undef:
727 case Intrinsic::spv_value_md:
728 case Intrinsic::spv_workgroup_size:
740 case SPIRV::OpTypeVoid:
741 case SPIRV::OpTypeBool:
742 case SPIRV::OpTypeInt:
743 case SPIRV::OpTypeFloat:
744 case SPIRV::OpTypeVector:
745 case SPIRV::OpTypeMatrix:
746 case SPIRV::OpTypeImage:
747 case SPIRV::OpTypeSampler:
748 case SPIRV::OpTypeSampledImage:
749 case SPIRV::OpTypeArray:
750 case SPIRV::OpTypeRuntimeArray:
751 case SPIRV::OpTypeStruct:
752 case SPIRV::OpTypeOpaque:
753 case SPIRV::OpTypePointer:
754 case SPIRV::OpTypeFunction:
755 case SPIRV::OpTypeEvent:
756 case SPIRV::OpTypeDeviceEvent:
757 case SPIRV::OpTypeReserveId:
758 case SPIRV::OpTypeQueue:
759 case SPIRV::OpTypePipe:
760 case SPIRV::OpTypeForwardPointer:
761 case SPIRV::OpTypePipeStorage:
762 case SPIRV::OpTypeNamedBarrier:
763 case SPIRV::OpTypeAccelerationStructureNV:
764 case SPIRV::OpTypeCooperativeMatrixNV:
765 case SPIRV::OpTypeCooperativeMatrixKHR:
775 if (
MI.getNumDefs() == 0)
778 for (
const auto &MO :
MI.all_defs()) {
780 if (
Reg.isPhysical()) {
785 if (
UseMI.getOpcode() != SPIRV::OpName) {
792 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
793 MI.isLifetimeMarker()) {
796 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
807 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
808 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
811 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
816 if (
MI.mayStore() ||
MI.isCall() ||
817 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
818 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
819 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
830 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
837void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
839 for (
const auto &MO :
MI.all_defs()) {
843 SmallVector<MachineInstr *, 4> UselessOpNames;
846 "There is still a use of the dead function.");
849 for (MachineInstr *OpNameMI : UselessOpNames) {
851 OpNameMI->eraseFromParent();
856void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
859 removeOpNamesForDeadMI(
MI);
860 MI.eraseFromParent();
863bool SPIRVInstructionSelector::select(MachineInstr &
I) {
864 resetVRegsType(*
I.getParent()->getParent());
866 assert(
I.getParent() &&
"Instruction should be in a basic block!");
867 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
872 removeDeadInstruction(
I);
879 if (Opcode == SPIRV::ASSIGN_TYPE) {
880 Register DstReg =
I.getOperand(0).getReg();
881 Register SrcReg =
I.getOperand(1).getReg();
884 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
885 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
886 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
887 Register SelectDstReg =
Def->getOperand(0).getReg();
888 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
890 assert(SuccessToSelectSelect);
892 Def->eraseFromParent();
899 bool Res = selectImpl(
I, *CoverageInfo);
901 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
902 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
906 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
918 }
else if (
I.getNumDefs() == 1) {
930 removeDeadInstruction(
I);
935 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
936 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
942 bool HasDefs =
I.getNumDefs() > 0;
945 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
946 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
947 if (spvSelect(ResVReg, ResType,
I)) {
949 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
960 case TargetOpcode::G_CONSTANT:
961 case TargetOpcode::G_FCONSTANT:
968 MachineInstr &
I)
const {
971 if (DstRC != SrcRC && SrcRC)
973 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
980bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
981 SPIRVTypeInst ResType,
982 MachineInstr &
I)
const {
983 const unsigned Opcode =
I.getOpcode();
985 return selectImpl(
I, *CoverageInfo);
987 case TargetOpcode::G_CONSTANT:
988 case TargetOpcode::G_FCONSTANT:
989 return selectConst(ResVReg, ResType,
I);
990 case TargetOpcode::G_GLOBAL_VALUE:
991 return selectGlobalValue(ResVReg,
I);
992 case TargetOpcode::G_IMPLICIT_DEF:
993 return selectOpUndef(ResVReg, ResType,
I);
994 case TargetOpcode::G_FREEZE:
995 return selectFreeze(ResVReg, ResType,
I);
997 case TargetOpcode::G_INTRINSIC:
998 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
999 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1000 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1001 return selectIntrinsic(ResVReg, ResType,
I);
1002 case TargetOpcode::G_BITREVERSE:
1003 return selectBitreverse(ResVReg, ResType,
I);
1005 case TargetOpcode::G_BUILD_VECTOR:
1006 return selectBuildVector(ResVReg, ResType,
I);
1007 case TargetOpcode::G_SPLAT_VECTOR:
1008 return selectSplatVector(ResVReg, ResType,
I);
1010 case TargetOpcode::G_SHUFFLE_VECTOR: {
1011 MachineBasicBlock &BB = *
I.getParent();
1012 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1015 .
addUse(
I.getOperand(1).getReg())
1016 .
addUse(
I.getOperand(2).getReg());
1017 for (
auto V :
I.getOperand(3).getShuffleMask())
1022 case TargetOpcode::G_MEMMOVE:
1023 case TargetOpcode::G_MEMCPY:
1024 case TargetOpcode::G_MEMSET:
1025 return selectMemOperation(ResVReg,
I);
1027 case TargetOpcode::G_ICMP:
1028 return selectICmp(ResVReg, ResType,
I);
1029 case TargetOpcode::G_FCMP:
1030 return selectFCmp(ResVReg, ResType,
I);
1032 case TargetOpcode::G_FRAME_INDEX:
1033 return selectFrameIndex(ResVReg, ResType,
I);
1035 case TargetOpcode::G_LOAD:
1036 return selectLoad(ResVReg, ResType,
I);
1037 case TargetOpcode::G_STORE:
1038 return selectStore(
I);
1040 case TargetOpcode::G_BR:
1041 return selectBranch(
I);
1042 case TargetOpcode::G_BRCOND:
1043 return selectBranchCond(
I);
1045 case TargetOpcode::G_PHI:
1046 return selectPhi(ResVReg,
I);
1048 case TargetOpcode::G_FPTOSI:
1049 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1050 case TargetOpcode::G_FPTOUI:
1051 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1053 case TargetOpcode::G_FPTOSI_SAT:
1054 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1055 case TargetOpcode::G_FPTOUI_SAT:
1056 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1058 case TargetOpcode::G_SITOFP:
1059 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1060 case TargetOpcode::G_UITOFP:
1061 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1063 case TargetOpcode::G_CTPOP:
1064 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1065 case TargetOpcode::G_SMIN:
1066 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1067 case TargetOpcode::G_UMIN:
1068 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1070 case TargetOpcode::G_SMAX:
1071 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1072 case TargetOpcode::G_UMAX:
1073 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1075 case TargetOpcode::G_SCMP:
1076 return selectSUCmp(ResVReg, ResType,
I,
true);
1077 case TargetOpcode::G_UCMP:
1078 return selectSUCmp(ResVReg, ResType,
I,
false);
1079 case TargetOpcode::G_LROUND:
1080 case TargetOpcode::G_LLROUND: {
1083 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1085 regForLround, *(
I.getParent()->getParent()));
1087 CL::round, GL::Round,
false);
1089 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1096 case TargetOpcode::G_STRICT_FMA:
1097 case TargetOpcode::G_FMA: {
1100 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1103 .
addUse(
I.getOperand(1).getReg())
1104 .
addUse(
I.getOperand(2).getReg())
1105 .
addUse(
I.getOperand(3).getReg())
1110 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1113 case TargetOpcode::G_STRICT_FLDEXP:
1114 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1116 case TargetOpcode::G_FPOW:
1117 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1118 case TargetOpcode::G_FPOWI:
1119 return selectFpowi(ResVReg, ResType,
I);
1121 case TargetOpcode::G_FEXP:
1122 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1123 case TargetOpcode::G_FEXP2:
1124 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1125 case TargetOpcode::G_FEXP10:
1126 return selectExp10(ResVReg, ResType,
I);
1128 case TargetOpcode::G_FMODF:
1129 return selectModf(ResVReg, ResType,
I);
1130 case TargetOpcode::G_FSINCOS:
1131 return selectSincos(ResVReg, ResType,
I);
1133 case TargetOpcode::G_FLOG:
1134 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1135 case TargetOpcode::G_FLOG2:
1136 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1137 case TargetOpcode::G_FLOG10:
1138 return selectLog10(ResVReg, ResType,
I);
1140 case TargetOpcode::G_FABS:
1141 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1142 case TargetOpcode::G_ABS:
1143 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1145 case TargetOpcode::G_FMINNUM:
1146 case TargetOpcode::G_FMINIMUM:
1147 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1148 case TargetOpcode::G_FMAXNUM:
1149 case TargetOpcode::G_FMAXIMUM:
1150 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1152 case TargetOpcode::G_FCOPYSIGN:
1153 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1155 case TargetOpcode::G_FCEIL:
1156 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1157 case TargetOpcode::G_FFLOOR:
1158 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1160 case TargetOpcode::G_FCOS:
1161 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1162 case TargetOpcode::G_FSIN:
1163 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1164 case TargetOpcode::G_FTAN:
1165 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1166 case TargetOpcode::G_FACOS:
1167 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1168 case TargetOpcode::G_FASIN:
1169 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1170 case TargetOpcode::G_FATAN:
1171 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1172 case TargetOpcode::G_FATAN2:
1173 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1174 case TargetOpcode::G_FCOSH:
1175 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1176 case TargetOpcode::G_FSINH:
1177 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1178 case TargetOpcode::G_FTANH:
1179 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1181 case TargetOpcode::G_STRICT_FSQRT:
1182 case TargetOpcode::G_FSQRT:
1183 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1185 case TargetOpcode::G_CTTZ:
1186 case TargetOpcode::G_CTTZ_ZERO_POISON:
1187 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1188 case TargetOpcode::G_CTLZ:
1189 case TargetOpcode::G_CTLZ_ZERO_POISON:
1190 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1192 case TargetOpcode::G_INTRINSIC_ROUND:
1193 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1194 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1195 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1196 case TargetOpcode::G_INTRINSIC_TRUNC:
1197 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1198 case TargetOpcode::G_FRINT:
1199 case TargetOpcode::G_FNEARBYINT:
1200 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1202 case TargetOpcode::G_SMULH:
1203 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1204 case TargetOpcode::G_UMULH:
1205 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1207 case TargetOpcode::G_SADDSAT:
1208 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1209 case TargetOpcode::G_UADDSAT:
1210 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1211 case TargetOpcode::G_SSUBSAT:
1212 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1213 case TargetOpcode::G_USUBSAT:
1214 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1216 case TargetOpcode::G_FFREXP:
1217 return selectFrexp(ResVReg, ResType,
I);
1219 case TargetOpcode::G_UADDO:
1220 return selectOverflowArith(ResVReg, ResType,
I,
1221 ResType->
getOpcode() == SPIRV::OpTypeVector
1222 ? SPIRV::OpIAddCarryV
1223 : SPIRV::OpIAddCarryS);
1224 case TargetOpcode::G_USUBO:
1225 return selectOverflowArith(ResVReg, ResType,
I,
1226 ResType->
getOpcode() == SPIRV::OpTypeVector
1227 ? SPIRV::OpISubBorrowV
1228 : SPIRV::OpISubBorrowS);
1229 case TargetOpcode::G_UMULO:
1230 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1231 case TargetOpcode::G_SMULO:
1232 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1234 case TargetOpcode::G_SEXT:
1235 return selectExt(ResVReg, ResType,
I,
true);
1236 case TargetOpcode::G_ANYEXT:
1237 case TargetOpcode::G_ZEXT:
1238 return selectExt(ResVReg, ResType,
I,
false);
1239 case TargetOpcode::G_TRUNC:
1240 return selectTrunc(ResVReg, ResType,
I);
1241 case TargetOpcode::G_FPTRUNC:
1242 case TargetOpcode::G_FPEXT:
1243 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1245 case TargetOpcode::G_PTRTOINT:
1246 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1247 case TargetOpcode::G_INTTOPTR:
1248 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1249 case TargetOpcode::G_BITCAST:
1250 return selectBitcast(ResVReg, ResType,
I);
1251 case TargetOpcode::G_ADDRSPACE_CAST:
1252 return selectAddrSpaceCast(ResVReg, ResType,
I);
1253 case TargetOpcode::G_PTR_ADD: {
1255 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1259 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1260 (*II).getOpcode() == TargetOpcode::COPY ||
1261 (*II).getOpcode() == SPIRV::OpVariable) &&
1262 getImm(
I.getOperand(2), MRI));
1264 bool IsGVInit =
false;
1268 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1269 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1270 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1271 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1281 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1294 "incompatible result and operand types in a bitcast");
1296 MachineInstrBuilder MIB =
1297 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1304 : SPIRV::OpInBoundsPtrAccessChain))
1308 .
addUse(
I.getOperand(2).getReg())
1311 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1315 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1317 .
addUse(
I.getOperand(2).getReg())
1326 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1329 .
addImm(
static_cast<uint32_t
>(
1330 SPIRV::Opcode::InBoundsPtrAccessChain))
1333 .
addUse(
I.getOperand(2).getReg());
1338 case TargetOpcode::G_ATOMICRMW_OR:
1339 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1340 case TargetOpcode::G_ATOMICRMW_ADD:
1341 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1342 case TargetOpcode::G_ATOMICRMW_AND:
1343 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1344 case TargetOpcode::G_ATOMICRMW_MAX:
1345 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1346 case TargetOpcode::G_ATOMICRMW_MIN:
1347 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1348 case TargetOpcode::G_ATOMICRMW_SUB:
1349 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1350 case TargetOpcode::G_ATOMICRMW_XOR:
1351 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1352 case TargetOpcode::G_ATOMICRMW_UMAX:
1353 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1354 case TargetOpcode::G_ATOMICRMW_UMIN:
1355 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1356 case TargetOpcode::G_ATOMICRMW_XCHG:
1357 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1359 case TargetOpcode::G_ATOMICRMW_FADD:
1360 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1361 case TargetOpcode::G_ATOMICRMW_FSUB:
1363 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1364 ResType->
getOpcode() == SPIRV::OpTypeVector
1366 : SPIRV::OpFNegate);
1367 case TargetOpcode::G_ATOMICRMW_FMIN:
1368 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1369 case TargetOpcode::G_ATOMICRMW_FMAX:
1370 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1372 case TargetOpcode::G_FENCE:
1373 return selectFence(
I);
1375 case TargetOpcode::G_STACKSAVE:
1376 return selectStackSave(ResVReg, ResType,
I);
1377 case TargetOpcode::G_STACKRESTORE:
1378 return selectStackRestore(
I);
1380 case TargetOpcode::G_UNMERGE_VALUES:
1383 case TargetOpcode::G_TRAP:
1384 case TargetOpcode::G_UBSANTRAP:
1385 return selectTrap(
I);
1390 case TargetOpcode::DBG_LABEL:
1392 case TargetOpcode::G_DEBUGTRAP:
1393 return selectDebugTrap(ResVReg, ResType,
I);
1400bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1401 SPIRVTypeInst ResType,
1402 MachineInstr &
I)
const {
1403 unsigned Opcode = SPIRV::OpNop;
1410bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1411 SPIRVTypeInst ResType,
1413 GL::GLSLExtInst GLInst,
1414 bool setMIFlags,
bool useMISrc,
1417 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1418 std::string DiagMsg;
1419 raw_string_ostream OS(DiagMsg);
1420 I.print(OS,
true,
false,
false,
false);
1421 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1424 return selectExtInst(ResVReg, ResType,
I,
1425 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1426 setMIFlags, useMISrc, SrcRegs);
1429bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1430 SPIRVTypeInst ResType,
1432 CL::OpenCLExtInst CLInst,
1433 bool setMIFlags,
bool useMISrc,
1435 return selectExtInst(ResVReg, ResType,
I,
1436 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1437 setMIFlags, useMISrc, SrcRegs);
1440bool SPIRVInstructionSelector::selectExtInst(
1441 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1442 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1444 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1445 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1446 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1450bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1451 SPIRVTypeInst ResType,
1454 bool setMIFlags,
bool useMISrc,
1457 for (
const auto &[InstructionSet, Opcode] : Insts) {
1461 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1464 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1469 const unsigned NumOps =
I.getNumOperands();
1472 I.getOperand(Index).getType() ==
1473 MachineOperand::MachineOperandType::MO_IntrinsicID)
1476 MIB.
add(
I.getOperand(Index));
1488bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1489 SPIRVTypeInst ResType,
1490 MachineInstr &
I)
const {
1491 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1492 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1493 for (
const auto &Ex : ExtInsts) {
1494 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1495 uint32_t Opcode = Ex.second;
1499 MachineIRBuilder MIRBuilder(
I);
1502 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1507 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1510 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1513 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1516 .
addImm(
static_cast<uint32_t
>(Ex.first))
1518 .
add(
I.getOperand(2))
1522 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1523 .
addDef(
I.getOperand(1).getReg())
1532bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1533 SPIRVTypeInst ResType,
1534 MachineInstr &
I)
const {
1535 Register CosResVReg =
I.getOperand(1).getReg();
1536 unsigned SrcIdx =
I.getNumExplicitDefs();
1541 MachineIRBuilder MIRBuilder(
I);
1543 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1548 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1551 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1553 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1556 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1558 .
add(
I.getOperand(SrcIdx))
1561 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1569 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1572 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1574 .
add(
I.getOperand(SrcIdx))
1576 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1579 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1581 .
add(
I.getOperand(SrcIdx))
1588bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1589 SPIRVTypeInst ResType,
1591 std::vector<Register> Srcs,
1592 unsigned Opcode)
const {
1593 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1603std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1604 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1605 SPIRVTypeInst I32Type)
const {
1608 if (ComponentCount == 1) {
1611 Parts.IsScalar =
true;
1612 Parts.Type = I32Type;
1620 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1621 SPIRV::OpVectorExtractDynamic))
1622 return std::nullopt;
1624 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1625 SPIRV::OpVectorExtractDynamic))
1626 return std::nullopt;
1630 MachineIRBuilder MIRBuilder(
I);
1631 Parts.IsScalar =
false;
1638 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1639 TII.get(SPIRV::OpVectorShuffle))
1644 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1649 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1650 TII.get(SPIRV::OpVectorShuffle))
1655 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1663bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1664 SPIRVTypeInst ResType,
1667 unsigned Opcode)
const {
1668 Register OpReg =
I.getOperand(1).getReg();
1671 MachineIRBuilder MIRBuilder(
I);
1673 SPIRVTypeInst I32VectorType =
1676 bool IsVector = NumElems > 1;
1677 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1680 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1684 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1687 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1690bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1691 SPIRVTypeInst ResType,
1694 unsigned Opcode)
const {
1695 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1698bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1699 SPIRVTypeInst ResType,
1702 unsigned Opcode)
const {
1704 if (ComponentCount > 2)
1705 return handle64BitOverflow(
1706 ResVReg, ResType,
I, SrcReg, Opcode,
1708 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1710 MachineIRBuilder MIRBuilder(
I);
1715 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1719 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1724 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1728 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1731 SplitParts &Parts = *MaybeParts;
1734 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1736 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1741 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1742 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1745bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1746 SPIRVTypeInst ResType,
1748 unsigned Opcode)
const {
1753 if (!STI.getTargetTriple().isVulkanOS())
1754 return selectUnOp(ResVReg, ResType,
I, Opcode);
1756 Register OpReg =
I.getOperand(1).getReg();
1759 : SPIRV::OpUConvert;
1763 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1765 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1767 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1773bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1774 SPIRVTypeInst ResType,
1776 unsigned Opcode)
const {
1778 Register SrcReg =
I.getOperand(1).getReg();
1783 unsigned DefOpCode = DefIt->getOpcode();
1784 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1787 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1788 DefOpCode = VRD->getOpcode();
1790 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1791 DefOpCode == TargetOpcode::G_CONSTANT ||
1792 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1798 uint32_t SpecOpcode = 0;
1800 case SPIRV::OpConvertPtrToU:
1801 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1803 case SPIRV::OpConvertUToPtr:
1804 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1809 TII.get(SPIRV::OpSpecConstantOp))
1819 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1823bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1824 SPIRVTypeInst ResType,
1825 MachineInstr &
I)
const {
1826 Register OpReg =
I.getOperand(1).getReg();
1827 SPIRVTypeInst OpType =
1831 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1841 if (
MemOp->isVolatile())
1842 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1843 if (
MemOp->isNonTemporal())
1844 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1846 if (!ST->isShader() &&
MemOp->getAlign().value())
1847 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1851 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1852 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1856 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1858 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1862 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1866 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1868 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1880 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1882 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1884 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1888bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1889 SPIRVTypeInst ResType,
1890 MachineInstr &
I)
const {
1892 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1897 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1898 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1900 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1902 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1906 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1910 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1911 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1912 I.getDebugLoc(),
I);
1916 MachineIRBuilder MIRBuilder(
I);
1918 if (
I.getNumMemOperands()) {
1919 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1920 if (MemOp->isAtomic())
1921 return selectAtomicLoad(ResVReg, ResType,
I);
1924 auto MIB = MIRBuilder.buildInstr(SPIRV::OpLoad)
1928 if (!
I.getNumMemOperands()) {
1929 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1931 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1940bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
1941 SPIRVTypeInst ResType,
1942 MachineInstr &
I)
const {
1943 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
1946 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1949 return diagnoseUnsupported(
I,
1950 "Lowering to SPIR-V of atomic load is only "
1951 "allowed for integer or floating point types");
1953 assert(
I.getNumMemOperands());
1954 const MachineMemOperand &MemOp = **
I.memoperands_begin();
1955 assert(MemOp.isAtomic());
1959 Register ScopeReg = buildI32Constant(Scope,
I);
1965 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
1966 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
1969 MachineIRBuilder MIRBuilder(
I);
1970 auto AtomicLoad = MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
1976 AtomicLoad.constrainAllUses(
TII,
TRI, RBI);
1980bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1982 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1983 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1988 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1989 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1991 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1996 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2000 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2001 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2002 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2003 TII.get(SPIRV::OpImageWrite))
2009 if (sampledTypeIsSignedInteger(LLVMHandleType))
2012 BMI.constrainAllUses(
TII,
TRI, RBI);
2017 if (
I.getNumMemOperands()) {
2018 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2019 if (MemOp->isAtomic())
2020 return selectAtomicStore(
I);
2023 MachineIRBuilder MIRBuilder(
I);
2024 auto MIB = MIRBuilder.buildInstr(SPIRV::OpStore).
addUse(Ptr).
addUse(StoreVal);
2025 if (!
I.getNumMemOperands()) {
2026 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2028 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2037bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2038 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2041 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2042 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2047 return diagnoseUnsupported(
I,
2048 "Lowering to SPIR-V of atomic store is only "
2049 "allowed for integer or floating point types");
2051 assert(
I.getNumMemOperands());
2052 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2053 assert(MemOp.isAtomic());
2057 Register ScopeReg = buildI32Constant(Scope,
I);
2063 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2064 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2067 MachineIRBuilder MIRBuilder(
I);
2068 auto AtomicStore = MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
2073 AtomicStore.constrainAllUses(
TII,
TRI, RBI);
2077bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2078 SPIRVTypeInst ResType,
2079 MachineInstr &
I)
const {
2080 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2088 const Register PtrsReg =
I.getOperand(2).getReg();
2089 const uint32_t Alignment =
I.getOperand(3).getImm();
2090 const Register MaskReg =
I.getOperand(4).getReg();
2091 const Register PassthruReg =
I.getOperand(5).getReg();
2092 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2096 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2107bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2108 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2115 const Register ValuesReg =
I.getOperand(1).getReg();
2116 const Register PtrsReg =
I.getOperand(2).getReg();
2117 const uint32_t Alignment =
I.getOperand(3).getImm();
2118 const Register MaskReg =
I.getOperand(4).getReg();
2119 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2123 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2132bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2133 const Twine &Msg)
const {
2134 const Function &
F =
I.getMF()->getFunction();
2135 F.getContext().diagnose(
2136 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
2140bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2141 SPIRVTypeInst ResType,
2142 MachineInstr &
I)
const {
2143 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2145 "llvm.stacksave intrinsic: this instruction requires the following "
2146 "SPIR-V extension: SPV_INTEL_variable_length_array",
2149 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2156bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2157 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2159 "llvm.stackrestore intrinsic: this instruction requires the following "
2160 "SPIR-V extension: SPV_INTEL_variable_length_array",
2162 if (!
I.getOperand(0).isReg())
2165 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2166 .
addUse(
I.getOperand(0).getReg())
2172SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2173 MachineIRBuilder MIRBuilder(
I);
2174 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2181 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2185 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2186 Type *ArrTy = ArrayType::get(ValTy, Num);
2188 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2191 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2198 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
2201 .
addImm(SPIRV::StorageClass::UniformConstant)
2212bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2215 Register DstReg =
I.getOperand(0).getReg();
2220 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2225 "Unable to determine pointee type size for OpCopyMemory");
2226 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2227 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2229 "OpCopyMemory requires the size to match the pointee type size");
2230 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2233 if (
I.getNumMemOperands()) {
2234 MachineIRBuilder MIRBuilder(
I);
2241bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2244 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2245 .
addUse(
I.getOperand(0).getReg())
2247 .
addUse(
I.getOperand(2).getReg());
2248 if (
I.getNumMemOperands()) {
2249 MachineIRBuilder MIRBuilder(
I);
2256bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2257 MachineInstr &
I)
const {
2258 Register SrcReg =
I.getOperand(1).getReg();
2259 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
2260 Register VarReg = getOrCreateMemSetGlobal(
I);
2263 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2265 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2267 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2271 if (!selectCopyMemory(
I, SrcReg))
2274 if (!selectCopyMemorySized(
I, SrcReg))
2277 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2278 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2283bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2284 SPIRVTypeInst ResType,
2287 unsigned NegateOpcode)
const {
2289 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2292 Register ScopeReg = buildI32Constant(Scope,
I);
2294 Register Ptr =
I.getOperand(1).getReg();
2295 uint32_t ScSem =
static_cast<uint32_t
>(
2299 Register MemSemReg = buildI32Constant(MemSem,
I);
2301 Register ValueReg =
I.getOperand(2).getReg();
2302 if (NegateOpcode != 0) {
2305 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2310 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2321bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2322 unsigned ArgI =
I.getNumOperands() - 1;
2324 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2325 SPIRVTypeInst SrcType =
2327 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2329 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2333 unsigned CurrentIndex = 0;
2334 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2335 Register ResVReg =
I.getOperand(i).getReg();
2338 LLT ResLLT = MRI->
getType(ResVReg);
2344 ResType = ScalarType;
2350 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2353 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2359 for (
unsigned j = 0;
j < NumElements; ++
j) {
2360 MIB.
addImm(CurrentIndex + j);
2362 CurrentIndex += NumElements;
2366 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2378bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2381 Register MemSemReg = buildI32Constant(MemSem,
I);
2383 uint32_t
Scope =
static_cast<uint32_t
>(
2385 Register ScopeReg = buildI32Constant(Scope,
I);
2387 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2394bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2395 SPIRVTypeInst ResType,
2397 unsigned Opcode)
const {
2398 Type *ResTy =
nullptr;
2402 "Not enough info to select the arithmetic with overflow instruction");
2405 "with overflow instruction");
2411 MachineIRBuilder MIRBuilder(
I);
2413 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2414 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2420 Register ZeroReg = buildZerosVal(ResType,
I);
2425 if (ResName.
size() > 0)
2430 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2433 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2434 MIB.
addUse(
I.getOperand(i).getReg());
2439 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2440 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2442 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2443 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2450 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2451 .
addDef(
I.getOperand(1).getReg())
2459bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2460 SPIRVTypeInst ResType,
2461 MachineInstr &
I)
const {
2463 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2464 Register Ptr =
I.getOperand(2).getReg();
2465 Register ScopeReg =
I.getOperand(5).getReg();
2466 Register MemSemEqReg =
I.getOperand(6).getReg();
2467 Register MemSemNeqReg =
I.getOperand(7).getReg();
2469 Register Val =
I.getOperand(4).getReg();
2473 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2492 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2499 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2511 case SPIRV::StorageClass::DeviceOnlyINTEL:
2512 case SPIRV::StorageClass::HostOnlyINTEL:
2521 bool IsGRef =
false;
2522 bool IsAllowedRefs =
2524 unsigned Opcode = It.getOpcode();
2525 if (Opcode == SPIRV::OpConstantComposite ||
2526 Opcode == SPIRV::OpSpecConstantComposite ||
2527 Opcode == SPIRV::OpVariable ||
2528 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2529 return IsGRef = true;
2530 return Opcode == SPIRV::OpName;
2532 return IsAllowedRefs && IsGRef;
2535Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2536 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2538 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2542SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2544 uint32_t Opcode)
const {
2545 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2546 TII.get(SPIRV::OpSpecConstantOp))
2554SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2555 SPIRVTypeInst SrcPtrTy)
const {
2556 SPIRVTypeInst GenericPtrTy =
2560 SPIRV::StorageClass::Generic),
2562 MachineFunction *MF =
I.getParent()->getParent();
2564 MachineInstrBuilder MIB = buildSpecConstantOp(
2566 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2576bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2577 SPIRVTypeInst ResType,
2578 MachineInstr &
I)
const {
2582 Register SrcPtr =
I.getOperand(1).getReg();
2586 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2587 ResType->
getOpcode() != SPIRV::OpTypePointer)
2588 return BuildCOPY(ResVReg, SrcPtr,
I);
2598 unsigned SpecOpcode =
2600 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2603 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2610 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2612 .constrainAllUses(
TII,
TRI, RBI);
2614 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2616 buildSpecConstantOp(
2618 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2619 .constrainAllUses(
TII,
TRI, RBI);
2626 return BuildCOPY(ResVReg, SrcPtr,
I);
2628 if ((SrcSC == SPIRV::StorageClass::Function &&
2629 DstSC == SPIRV::StorageClass::Private) ||
2630 (DstSC == SPIRV::StorageClass::Function &&
2631 SrcSC == SPIRV::StorageClass::Private))
2632 return BuildCOPY(ResVReg, SrcPtr,
I);
2636 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2639 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2642 SPIRVTypeInst GenericPtrTy =
2661 return selectUnOp(ResVReg, ResType,
I,
2662 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2664 return selectUnOp(ResVReg, ResType,
I,
2665 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2667 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2669 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2679 return SPIRV::OpFOrdEqual;
2681 return SPIRV::OpFOrdGreaterThanEqual;
2683 return SPIRV::OpFOrdGreaterThan;
2685 return SPIRV::OpFOrdLessThanEqual;
2687 return SPIRV::OpFOrdLessThan;
2689 return SPIRV::OpFOrdNotEqual;
2691 return SPIRV::OpOrdered;
2693 return SPIRV::OpFUnordEqual;
2695 return SPIRV::OpFUnordGreaterThanEqual;
2697 return SPIRV::OpFUnordGreaterThan;
2699 return SPIRV::OpFUnordLessThanEqual;
2701 return SPIRV::OpFUnordLessThan;
2703 return SPIRV::OpFUnordNotEqual;
2705 return SPIRV::OpUnordered;
2715 return SPIRV::OpIEqual;
2717 return SPIRV::OpINotEqual;
2719 return SPIRV::OpSGreaterThanEqual;
2721 return SPIRV::OpSGreaterThan;
2723 return SPIRV::OpSLessThanEqual;
2725 return SPIRV::OpSLessThan;
2727 return SPIRV::OpUGreaterThanEqual;
2729 return SPIRV::OpUGreaterThan;
2731 return SPIRV::OpULessThanEqual;
2733 return SPIRV::OpULessThan;
2742 return SPIRV::OpPtrEqual;
2744 return SPIRV::OpPtrNotEqual;
2755 return SPIRV::OpLogicalEqual;
2757 return SPIRV::OpLogicalNotEqual;
2791bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2792 SPIRVTypeInst ResType,
2794 unsigned OpAnyOrAll)
const {
2795 assert(
I.getNumOperands() == 3);
2796 assert(
I.getOperand(2).isReg());
2798 Register InputRegister =
I.getOperand(2).getReg();
2801 assert(InputType &&
"VReg has no type assigned");
2804 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2805 if (IsBoolTy && !IsVectorTy) {
2806 assert(ResVReg ==
I.getOperand(0).getReg());
2807 return BuildCOPY(ResVReg, InputRegister,
I);
2811 unsigned SpirvNotEqualId =
2812 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2814 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2819 IsBoolTy ? InputRegister
2827 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2829 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2846bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2847 SPIRVTypeInst ResType,
2848 MachineInstr &
I)
const {
2849 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2852bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2853 SPIRVTypeInst ResType,
2854 MachineInstr &
I)
const {
2855 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2859bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2860 SPIRVTypeInst ResType,
2861 MachineInstr &
I)
const {
2862 assert(
I.getNumOperands() == 4);
2863 assert(
I.getOperand(2).isReg());
2864 assert(
I.getOperand(3).isReg());
2866 [[maybe_unused]] SPIRVTypeInst VecType =
2871 "dot product requires a vector of at least 2 components");
2873 [[maybe_unused]] SPIRVTypeInst EltType =
2882 .
addUse(
I.getOperand(2).getReg())
2883 .
addUse(
I.getOperand(3).getReg())
2888bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2889 SPIRVTypeInst ResType,
2892 assert(
I.getNumOperands() == 4);
2893 assert(
I.getOperand(2).isReg());
2894 assert(
I.getOperand(3).isReg());
2897 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2901 .
addUse(
I.getOperand(2).getReg())
2902 .
addUse(
I.getOperand(3).getReg())
2909bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2910 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2911 assert(
I.getNumOperands() == 4);
2912 assert(
I.getOperand(2).isReg());
2913 assert(
I.getOperand(3).isReg());
2917 Register Vec0 =
I.getOperand(2).getReg();
2918 Register Vec1 =
I.getOperand(3).getReg();
2922 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2931 "dot product requires a vector of at least 2 components");
2934 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2944 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2955 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2967bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2968 SPIRVTypeInst ResType,
2969 MachineInstr &
I)
const {
2971 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2974 .
addUse(
I.getOperand(2).getReg())
2979bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2980 SPIRVTypeInst ResType,
2981 MachineInstr &
I)
const {
2983 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2986 .
addUse(
I.getOperand(2).getReg())
2991template <
bool Signed>
2992bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2993 SPIRVTypeInst ResType,
2994 MachineInstr &
I)
const {
2995 assert(
I.getNumOperands() == 5);
2996 assert(
I.getOperand(2).isReg());
2997 assert(
I.getOperand(3).isReg());
2998 assert(
I.getOperand(4).isReg());
3001 Register Acc =
I.getOperand(2).getReg();
3005 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3007 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3012 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3015 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3027template <
bool Signed>
3028bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3029 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3030 assert(
I.getNumOperands() == 5);
3031 assert(
I.getOperand(2).isReg());
3032 assert(
I.getOperand(3).isReg());
3033 assert(
I.getOperand(4).isReg());
3036 Register Acc =
I.getOperand(2).getReg();
3042 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3046 for (
unsigned i = 0; i < 4; i++) {
3069 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3089 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3104bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3105 SPIRVTypeInst ResType,
3106 MachineInstr &
I)
const {
3107 assert(
I.getNumOperands() == 3);
3108 assert(
I.getOperand(2).isReg());
3110 Register VZero = buildZerosValF(ResType,
I);
3111 Register VOne = buildOnesValF(ResType,
I);
3113 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3116 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3118 .
addUse(
I.getOperand(2).getReg())
3125bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3126 SPIRVTypeInst ResType,
3127 MachineInstr &
I)
const {
3128 assert(
I.getNumOperands() == 3);
3129 assert(
I.getOperand(2).isReg());
3131 Register InputRegister =
I.getOperand(2).getReg();
3133 auto &
DL =
I.getDebugLoc();
3143 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3145 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3153 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3158 if (NeedsConversion) {
3159 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3170bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3171 SPIRVTypeInst ResType,
3173 unsigned Opcode)
const {
3177 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3183 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3184 BMI.addUse(
I.getOperand(J).getReg());
3191bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3194 bool WithGroupSync)
const {
3196 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3198 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3200 assert(((Scope != SPIRV::Scope::Workgroup) ||
3201 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3202 "Workgroup Scope must set WorkGroupMemory semantic "
3203 "in Barrier instruction");
3205 assert(((Scope != SPIRV::Scope::Device) ||
3206 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3207 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3208 "Device Scope must set UniformMemory and ImageMemory semantic "
3209 "in Barrier instruction");
3215 if (WithGroupSync) {
3216 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3220 Register ScopeReg = buildI32Constant(Scope,
I);
3221 Register MemSemReg = buildI32Constant(MemSem,
I);
3223 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3227bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3228 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3233 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3234 SPIRV::OpGroupNonUniformBallot))
3239 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3244 .
addImm(SPIRV::GroupOperation::Reduce)
3251bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3252 SPIRVTypeInst ResType,
3253 MachineInstr &
I)
const {
3258 Register InputReg =
I.getOperand(2).getReg();
3263 bool IsVector = NumElems > 1;
3276 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3277 SPIRV::OpGroupNonUniformAllEqual);
3282 ElementResults.
reserve(NumElems);
3284 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3297 ElemInput = Extracted;
3303 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3314 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3325bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3326 SPIRVTypeInst ResType,
3327 MachineInstr &
I)
const {
3329 assert(
I.getNumOperands() == 3);
3331 auto Op =
I.getOperand(2);
3343 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3365 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3369 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3376bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3377 SPIRVTypeInst ResType,
3379 bool IsUnsigned)
const {
3380 return selectWaveReduce(
3381 ResVReg, ResType,
I, IsUnsigned,
3382 [&](
Register InputRegister,
bool IsUnsigned) {
3383 const bool IsFloatTy =
3385 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3386 : SPIRV::OpGroupNonUniformSMax;
3387 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3391bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3392 SPIRVTypeInst ResType,
3394 bool IsUnsigned)
const {
3395 return selectWaveReduce(
3396 ResVReg, ResType,
I, IsUnsigned,
3397 [&](
Register InputRegister,
bool IsUnsigned) {
3398 const bool IsFloatTy =
3400 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3401 : SPIRV::OpGroupNonUniformSMin;
3402 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3406bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3407 SPIRVTypeInst ResType,
3408 MachineInstr &
I)
const {
3409 return selectWaveReduce(ResVReg, ResType,
I,
false,
3410 [&](
Register InputRegister,
bool IsUnsigned) {
3412 InputRegister, SPIRV::OpTypeFloat);
3413 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3414 : SPIRV::OpGroupNonUniformIAdd;
3418bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3419 SPIRVTypeInst ResType,
3420 MachineInstr &
I)
const {
3421 return selectWaveReduce(ResVReg, ResType,
I,
false,
3422 [&](
Register InputRegister,
bool IsUnsigned) {
3424 InputRegister, SPIRV::OpTypeFloat);
3425 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3426 : SPIRV::OpGroupNonUniformIMul;
3430template <
typename PickOpcodeFn>
3431bool SPIRVInstructionSelector::selectWaveReduce(
3432 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3433 PickOpcodeFn &&PickOpcode)
const {
3434 assert(
I.getNumOperands() == 3);
3435 assert(
I.getOperand(2).isReg());
3437 Register InputRegister =
I.getOperand(2).getReg();
3444 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3450 .
addImm(SPIRV::GroupOperation::Reduce)
3451 .
addUse(
I.getOperand(2).getReg())
3456bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3457 SPIRVTypeInst ResType,
3459 unsigned Opcode)
const {
3460 return selectWaveReduce(
3461 ResVReg, ResType,
I,
false,
3462 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3465bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3466 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3467 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3468 [&](
Register InputRegister,
bool IsUnsigned) {
3470 InputRegister, SPIRV::OpTypeFloat);
3472 ? SPIRV::OpGroupNonUniformFAdd
3473 : SPIRV::OpGroupNonUniformIAdd;
3477bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3478 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3479 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3480 [&](
Register InputRegister,
bool IsUnsigned) {
3482 InputRegister, SPIRV::OpTypeFloat);
3484 ? SPIRV::OpGroupNonUniformFMul
3485 : SPIRV::OpGroupNonUniformIMul;
3489template <
typename PickOpcodeFn>
3490bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3491 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3492 PickOpcodeFn &&PickOpcode)
const {
3493 assert(
I.getNumOperands() == 3);
3494 assert(
I.getOperand(2).isReg());
3496 Register InputRegister =
I.getOperand(2).getReg();
3503 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3509 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3510 .
addUse(
I.getOperand(2).getReg())
3515bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3516 SPIRVTypeInst ResType,
3519 assert(
I.getNumOperands() == 3);
3520 assert(
I.getOperand(2).isReg());
3522 Register InputRegister =
I.getOperand(2).getReg();
3528 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3539bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3540 SPIRVTypeInst ResType,
3545 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3550 : SPIRV::OpUConvert;
3554 ShiftOp = SPIRV::OpShiftRightLogicalV;
3559 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3560 TII.get(SPIRV::OpConstantComposite))
3563 for (
unsigned It = 0; It <
N; ++It)
3567 ShiftConst = CompositeReg;
3572 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3577 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3582 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3587 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3590bool SPIRVInstructionSelector::handle64BitOverflow(
3592 unsigned int Opcode,
3599 "handle64BitOverflow should only be used for integer types");
3601 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
3603 MachineIRBuilder MIRBuilder(
I);
3605 SPIRVTypeInst I64x2Type =
3607 SPIRVTypeInst Vec2ResType =
3610 std::vector<Register> PartialRegs;
3612 unsigned CurrentComponent = 0;
3613 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3617 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3618 TII.get(SPIRV::OpVectorShuffle))
3623 .
addImm(CurrentComponent)
3624 .
addImm(CurrentComponent + 1);
3634 PartialRegs.push_back(SubVecReg);
3637 if (CurrentComponent != ComponentCount) {
3643 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
3644 SPIRV::OpVectorExtractDynamic))
3653 PartialRegs.push_back(FinalElemResReg);
3657 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
3658 SPIRV::OpCompositeConstruct);
3661bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
3662 SPIRVTypeInst ResType,
3666 if (ComponentCount > 2)
3667 return handle64BitOverflow(
3668 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
3670 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
3672 MachineIRBuilder MIRBuilder(
I);
3676 I32Type, 2 * ComponentCount, MIRBuilder,
false);
3680 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
3685 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
3689 auto MaybeParts = splitEvenOddLanes(Reverse32, ComponentCount,
I, I32Type);
3692 SplitParts &Parts = *MaybeParts;
3698 if (!selectOpWithSrcs(SwappedVec, VecI32Type,
I, {Parts.High, Parts.Low},
3699 SPIRV::OpCompositeConstruct))
3703 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
3706bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
3707 SPIRVTypeInst ResType,
3711 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3719bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3720 SPIRVTypeInst ResType,
3721 MachineInstr &
I)
const {
3722 Register OpReg =
I.getOperand(1).getReg();
3730 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3732 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3734 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
3736 return SPIRVInstructionSelector::diagnoseUnsupported(
3737 I,
"G_BITREVERSE only support 16,32,64 bits.");
3741 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3752 unsigned AndOp = SPIRV::OpBitwiseAndS;
3753 unsigned OrOp = SPIRV::OpBitwiseOrS;
3754 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
3755 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
3757 AndOp = SPIRV::OpBitwiseAndV;
3758 OrOp = SPIRV::OpBitwiseOrV;
3759 ShlOp = SPIRV::OpShiftLeftLogicalV;
3760 ShrOp = SPIRV::OpShiftRightLogicalV;
3766 const unsigned Shift) ->
Register {
3774 Register MaskReg = CreateConst(Mask);
3775 Register ShiftReg = CreateConst(Shift);
3782 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
3783 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
3784 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
3785 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
3786 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
3794 uint64_t
Mask = ~0ull;
3795 while ((Shift >>= 1) > 0) {
3802 return BuildCOPY(ResVReg, Result,
I);
3805bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3806 SPIRVTypeInst ResType,
3807 MachineInstr &
I)
const {
3813 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3815 Register OpReg =
I.getOperand(1).getReg();
3816 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
3817 if (
Def->getOpcode() == TargetOpcode::COPY)
3820 switch (
Def->getOpcode()) {
3821 case SPIRV::ASSIGN_TYPE:
3822 if (MachineInstr *AssignToDef =
3824 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3825 Reg =
Def->getOperand(2).getReg();
3828 case SPIRV::OpUndef:
3829 Reg =
Def->getOperand(1).getReg();
3832 unsigned DestOpCode;
3834 DestOpCode = SPIRV::OpConstantNull;
3836 DestOpCode = TargetOpcode::COPY;
3839 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3840 .
addDef(
I.getOperand(0).getReg())
3848bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3849 SPIRVTypeInst ResType,
3850 MachineInstr &
I)
const {
3852 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3854 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3858 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3863 for (
unsigned i =
I.getNumExplicitDefs();
3864 i <
I.getNumExplicitOperands() && IsConst; ++i)
3868 if (!IsConst &&
N < 2)
3870 "There must be at least two constituent operands in a vector");
3873 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3874 TII.get(IsConst ? SPIRV::OpConstantComposite
3875 : SPIRV::OpCompositeConstruct))
3878 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3879 MIB.
addUse(
I.getOperand(i).getReg());
3884bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3885 SPIRVTypeInst ResType,
3886 MachineInstr &
I)
const {
3888 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3890 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3896 if (!
I.getOperand(
OpIdx).isReg())
3903 if (!IsConst &&
N < 2)
3905 "There must be at least two constituent operands in a vector");
3908 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3909 TII.get(IsConst ? SPIRV::OpConstantComposite
3910 : SPIRV::OpCompositeConstruct))
3913 for (
unsigned i = 0; i <
N; ++i)
3919bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3920 SPIRVTypeInst ResType,
3921 MachineInstr &
I)
const {
3926 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3928 Opcode = SPIRV::OpDemoteToHelperInvocation;
3930 Opcode = SPIRV::OpKill;
3932 if (MachineInstr *NextI =
I.getNextNode()) {
3934 NextI->eraseFromParent();
3944bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3945 SPIRVTypeInst ResType,
unsigned CmpOpc,
3946 MachineInstr &
I)
const {
3947 Register Cmp0 =
I.getOperand(2).getReg();
3948 Register Cmp1 =
I.getOperand(3).getReg();
3951 "CMP operands should have the same type");
3952 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3962bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3963 SPIRVTypeInst ResType,
3964 MachineInstr &
I)
const {
3965 auto Pred =
I.getOperand(1).getPredicate();
3968 Register CmpOperand =
I.getOperand(2).getReg();
3973 Register Op1 =
I.getOperand(3).getReg();
3977 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
3982 I.getOperand(3).setReg(NewOp1);
3988 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3992SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3993 SPIRVTypeInst ResType)
const {
3995 SPIRVTypeInst SpvI32Ty =
3998 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4005 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4008 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4011 .
addImm(APInt(32, Val).getZExtValue());
4013 GR.
add(ConstInt,
MI);
4020Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4021 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4023 SPIRVTypeInst SpvI32Ty =
4025 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4030 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4031 MachineInstr *
MI =
nullptr;
4035 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4039 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4040 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4046 GR.
add(ConstInt,
MI);
4051bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4052 SPIRVTypeInst ResType,
4053 MachineInstr &
I)
const {
4055 return selectCmp(ResVReg, ResType, CmpOp,
I);
4058bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4059 SPIRVTypeInst ResType,
4060 MachineInstr &
I)
const {
4062 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4069 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4070 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4073 MachineIRBuilder MIRBuilder(
I);
4080 APFloat ConstVal(3.3219280948873623);
4084 APFloat::rmNearestTiesToEven, &LosesInfo);
4088 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4089 ? SPIRV::OpVectorTimesScalar
4092 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4093 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4095 if (!selectExtInst(ResVReg, ResType,
I,
4096 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4106Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4107 MachineInstr &
I)
const {
4110 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4115bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4121 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4129 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4132 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4133 Def->getOpcode() == SPIRV::OpConstantI)
4146 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4147 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4149 Intrinsic::spv_const_composite)) {
4150 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4151 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4152 if (!IsZero(
Def->getOperand(i).getReg()))
4161Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4162 MachineInstr &
I)
const {
4166 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4171Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4172 MachineInstr &
I)
const {
4176 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4182 SPIRVTypeInst ResType,
4183 MachineInstr &
I)
const {
4187 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4192bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4193 SPIRVTypeInst ResType,
4194 MachineInstr &
I)
const {
4195 Register SelectFirstArg =
I.getOperand(2).getReg();
4196 Register SelectSecondArg =
I.getOperand(3).getReg();
4205 SPIRV::OpTypeVector;
4212 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4213 }
else if (IsPtrTy) {
4214 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4216 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4219 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4220 "boolean condition");
4222 Opcode = SPIRV::OpSelectSFSCond;
4223 }
else if (IsPtrTy) {
4224 Opcode = SPIRV::OpSelectSPSCond;
4226 Opcode = SPIRV::OpSelectSISCond;
4229 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4232 .
addUse(
I.getOperand(1).getReg())
4241bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4242 SPIRVTypeInst ResType,
4244 MachineInstr &InsertAt,
4245 bool IsSigned)
const {
4247 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4248 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4249 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4251 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4263bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4264 SPIRVTypeInst ResType,
4265 MachineInstr &
I,
bool IsSigned,
4266 unsigned Opcode)
const {
4267 Register SrcReg =
I.getOperand(1).getReg();
4273 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4278 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
4280 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4283bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4284 SPIRVTypeInst ResType, MachineInstr &
I,
4285 bool IsSigned)
const {
4286 Register SrcReg =
I.getOperand(1).getReg();
4288 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4292 if (ResType == SrcType)
4293 return BuildCOPY(ResVReg, SrcReg,
I);
4295 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4296 return selectUnOp(ResVReg, ResType,
I, Opcode);
4299bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4300 SPIRVTypeInst ResType,
4302 bool IsSigned)
const {
4303 MachineIRBuilder MIRBuilder(
I);
4304 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4316 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4319 .
addUse(
I.getOperand(1).getReg())
4320 .
addUse(
I.getOperand(2).getReg())
4325 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4328 .
addUse(
I.getOperand(1).getReg())
4329 .
addUse(
I.getOperand(2).getReg())
4337 unsigned SelectOpcode =
4338 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4343 .
addUse(buildOnesVal(
true, ResType,
I))
4344 .
addUse(buildZerosVal(ResType,
I))
4351 .
addUse(buildOnesVal(
false, ResType,
I))
4356bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4359 SPIRVTypeInst IntTy,
4360 SPIRVTypeInst BoolTy)
const {
4363 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4364 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4366 Register One = buildOnesVal(
false, IntTy,
I);
4374 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4383bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4384 SPIRVTypeInst ResType,
4385 MachineInstr &
I)
const {
4386 Register IntReg =
I.getOperand(1).getReg();
4389 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4390 if (ArgType == ResType)
4391 return BuildCOPY(ResVReg, IntReg,
I);
4393 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4394 return selectUnOp(ResVReg, ResType,
I, Opcode);
4397bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4398 SPIRVTypeInst ResType,
4399 MachineInstr &
I)
const {
4400 unsigned Opcode =
I.getOpcode();
4401 unsigned TpOpcode = ResType->
getOpcode();
4403 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4404 assert(Opcode == TargetOpcode::G_CONSTANT &&
4405 I.getOperand(1).getCImm()->isZero());
4406 MachineBasicBlock &DepMBB =
I.getMF()->front();
4409 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4416 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4419bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4420 SPIRVTypeInst ResType,
4421 MachineInstr &
I)
const {
4422 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4429bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4430 SPIRVTypeInst ResType,
4431 MachineInstr &
I)
const {
4433 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4437 .
addUse(
I.getOperand(3).getReg())
4439 .
addUse(
I.getOperand(2).getReg());
4440 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4446bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4447 SPIRVTypeInst ResType,
4448 MachineInstr &
I)
const {
4449 Type *MaybeResTy =
nullptr;
4454 "Expected aggregate type for extractv instruction");
4456 SPIRV::AccessQualifier::ReadWrite,
false);
4460 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4463 .
addUse(
I.getOperand(2).getReg());
4464 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4470bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4471 SPIRVTypeInst ResType,
4472 MachineInstr &
I)
const {
4473 if (
getImm(
I.getOperand(4), MRI))
4474 return selectInsertVal(ResVReg, ResType,
I);
4476 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4479 .
addUse(
I.getOperand(2).getReg())
4480 .
addUse(
I.getOperand(3).getReg())
4481 .
addUse(
I.getOperand(4).getReg())
4486bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4487 SPIRVTypeInst ResType,
4488 MachineInstr &
I)
const {
4489 if (
getImm(
I.getOperand(3), MRI))
4490 return selectExtractVal(ResVReg, ResType,
I);
4492 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4495 .
addUse(
I.getOperand(2).getReg())
4496 .
addUse(
I.getOperand(3).getReg())
4501bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4502 SPIRVTypeInst ResType,
4503 MachineInstr &
I)
const {
4504 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4510 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4511 : SPIRV::OpAccessChain)
4512 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4513 :
SPIRV::OpPtrAccessChain);
4515 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4519 .
addUse(
I.getOperand(3).getReg());
4521 (Opcode == SPIRV::OpPtrAccessChain ||
4522 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4523 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4524 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4527 const unsigned StartingIndex =
4528 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4531 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4532 Res.addUse(
I.getOperand(i).getReg());
4533 Res.constrainAllUses(
TII,
TRI, RBI);
4538bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4540 unsigned Lim =
I.getNumExplicitOperands();
4541 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4542 Register OpReg =
I.getOperand(i).getReg();
4543 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4545 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4546 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4547 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4554 MachineFunction *MF =
I.getMF();
4566 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4567 TII.get(SPIRV::OpSpecConstantOp))
4570 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4572 GR.
add(OpDefine, MIB);
4578bool SPIRVInstructionSelector::selectDerivativeInst(
4579 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4580 const unsigned DPdOpCode)
const {
4583 errorIfInstrOutsideShader(
I);
4588 Register SrcReg =
I.getOperand(2).getReg();
4593 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4596 .
addUse(
I.getOperand(2).getReg());
4598 MachineIRBuilder MIRBuilder(
I);
4601 if (componentCount != 1)
4605 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4609 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4614 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4619 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4627bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4628 SPIRVTypeInst ResType,
4629 MachineInstr &
I)
const {
4633 case Intrinsic::spv_load:
4634 return selectLoad(ResVReg, ResType,
I);
4635 case Intrinsic::spv_atomic_load:
4636 return selectAtomicLoad(ResVReg, ResType,
I);
4637 case Intrinsic::spv_store:
4638 return selectStore(
I);
4639 case Intrinsic::spv_atomic_store:
4640 return selectAtomicStore(
I);
4641 case Intrinsic::spv_extractv:
4642 return selectExtractVal(ResVReg, ResType,
I);
4643 case Intrinsic::spv_insertv:
4644 return selectInsertVal(ResVReg, ResType,
I);
4645 case Intrinsic::spv_extractelt:
4646 return selectExtractElt(ResVReg, ResType,
I);
4647 case Intrinsic::spv_insertelt:
4648 return selectInsertElt(ResVReg, ResType,
I);
4649 case Intrinsic::spv_gep:
4650 return selectGEP(ResVReg, ResType,
I);
4651 case Intrinsic::spv_bitcast: {
4652 Register OpReg =
I.getOperand(2).getReg();
4653 SPIRVTypeInst OpType =
4657 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4659 case Intrinsic::spv_unref_global:
4660 case Intrinsic::spv_init_global: {
4661 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4666 Register GVarVReg =
MI->getOperand(0).getReg();
4667 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4672 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4674 MI->eraseFromParent();
4678 case Intrinsic::spv_undef: {
4679 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4685 case Intrinsic::spv_named_boolean_spec_constant: {
4686 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4687 : SPIRV::OpSpecConstantFalse;
4689 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4690 .
addDef(
I.getOperand(0).getReg())
4693 unsigned SpecId =
I.getOperand(2).getImm();
4695 SPIRV::Decoration::SpecId, {SpecId});
4699 case Intrinsic::spv_const_composite: {
4701 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4707 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4709 std::function<bool(
Register)> HasSpecConstOperand =
4719 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
4720 J < Def->getNumExplicitOperands(); ++J) {
4721 if (
Def->getOperand(J).isReg() &&
4722 HasSpecConstOperand(
Def->getOperand(J).getReg()))
4728 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
4729 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
4730 : SPIRV::OpConstantComposite;
4731 unsigned ContinuedOpc = HasSpecConst
4732 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
4733 : SPIRV::OpConstantCompositeContinuedINTEL;
4734 MachineIRBuilder MIR(
I);
4736 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
4738 for (
auto *Instr : Instructions) {
4739 Instr->setDebugLoc(
I.getDebugLoc());
4744 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4751 case Intrinsic::spv_assign_name: {
4752 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4753 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4754 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4755 i <
I.getNumExplicitOperands(); ++i) {
4756 MIB.
addImm(
I.getOperand(i).getImm());
4761 case Intrinsic::spv_switch: {
4762 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4763 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4764 if (
I.getOperand(i).isReg())
4765 MIB.
addReg(
I.getOperand(i).getReg());
4766 else if (
I.getOperand(i).isCImm())
4767 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4768 else if (
I.getOperand(i).isMBB())
4769 MIB.
addMBB(
I.getOperand(i).getMBB());
4776 case Intrinsic::spv_loop_merge: {
4777 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4778 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4779 if (
I.getOperand(i).isMBB())
4780 MIB.
addMBB(
I.getOperand(i).getMBB());
4787 case Intrinsic::spv_loop_control_intel: {
4789 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4790 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4795 case Intrinsic::spv_selection_merge: {
4797 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4798 assert(
I.getOperand(1).isMBB() &&
4799 "operand 1 to spv_selection_merge must be a basic block");
4800 MIB.
addMBB(
I.getOperand(1).getMBB());
4801 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4805 case Intrinsic::spv_cmpxchg:
4806 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4807 case Intrinsic::spv_unreachable:
4808 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4811 case Intrinsic::spv_abort:
4812 return selectAbort(
I);
4813 case Intrinsic::spv_alloca:
4814 return selectFrameIndex(ResVReg, ResType,
I);
4815 case Intrinsic::spv_alloca_array:
4816 return selectAllocaArray(ResVReg, ResType,
I);
4817 case Intrinsic::spv_assume:
4819 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4820 .
addUse(
I.getOperand(1).getReg())
4825 case Intrinsic::spv_expect:
4827 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4830 .
addUse(
I.getOperand(2).getReg())
4831 .
addUse(
I.getOperand(3).getReg())
4836 case Intrinsic::arithmetic_fence:
4837 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4838 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4841 .
addUse(
I.getOperand(2).getReg())
4845 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4847 case Intrinsic::spv_thread_id:
4853 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4855 case Intrinsic::spv_thread_id_in_group:
4861 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4863 case Intrinsic::spv_group_id:
4869 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4871 case Intrinsic::spv_flattened_thread_id_in_group:
4878 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4880 case Intrinsic::spv_workgroup_size:
4881 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4883 case Intrinsic::spv_global_size:
4884 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4886 case Intrinsic::spv_global_offset:
4887 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4889 case Intrinsic::spv_num_workgroups:
4890 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4892 case Intrinsic::spv_subgroup_size:
4893 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4895 case Intrinsic::spv_num_subgroups:
4896 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4898 case Intrinsic::spv_subgroup_id:
4899 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4900 case Intrinsic::spv_subgroup_local_invocation_id:
4901 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4902 ResVReg, ResType,
I);
4903 case Intrinsic::spv_subgroup_max_size:
4904 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4906 case Intrinsic::spv_fdot:
4907 return selectFloatDot(ResVReg, ResType,
I);
4908 case Intrinsic::spv_udot:
4909 case Intrinsic::spv_sdot:
4910 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4912 return selectIntegerDot(ResVReg, ResType,
I,
4913 IID == Intrinsic::spv_sdot);
4914 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4915 case Intrinsic::spv_dot4add_i8packed:
4916 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4918 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4919 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4920 case Intrinsic::spv_dot4add_u8packed:
4921 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4923 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4924 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4925 case Intrinsic::spv_all:
4926 return selectAll(ResVReg, ResType,
I);
4927 case Intrinsic::spv_any:
4928 return selectAny(ResVReg, ResType,
I);
4929 case Intrinsic::spv_cross:
4930 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4931 case Intrinsic::spv_distance:
4932 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4933 case Intrinsic::spv_lerp:
4934 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4935 case Intrinsic::spv_length:
4936 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4937 case Intrinsic::spv_degrees:
4938 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4939 case Intrinsic::spv_faceforward:
4940 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4941 case Intrinsic::spv_frac:
4942 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4943 case Intrinsic::spv_isinf:
4944 return selectOpIsInf(ResVReg, ResType,
I);
4945 case Intrinsic::spv_isnan:
4946 return selectOpIsNan(ResVReg, ResType,
I);
4947 case Intrinsic::spv_normalize:
4948 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4949 case Intrinsic::spv_refract:
4950 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4951 case Intrinsic::spv_reflect:
4952 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4953 case Intrinsic::spv_rsqrt:
4954 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4955 case Intrinsic::spv_sign:
4956 return selectSign(ResVReg, ResType,
I);
4957 case Intrinsic::spv_smoothstep:
4958 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4959 case Intrinsic::spv_firstbituhigh:
4960 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4961 case Intrinsic::spv_firstbitshigh:
4962 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4963 case Intrinsic::spv_firstbitlow:
4964 return selectFirstBitLow(ResVReg, ResType,
I);
4965 case Intrinsic::spv_all_memory_barrier:
4966 return selectBarrierInst(
I, SPIRV::Scope::Device,
4967 SPIRV::MemorySemantics::UniformMemory |
4968 SPIRV::MemorySemantics::ImageMemory |
4969 SPIRV::MemorySemantics::WorkgroupMemory,
4971 case Intrinsic::spv_all_memory_barrier_with_group_sync:
4972 return selectBarrierInst(
I, SPIRV::Scope::Device,
4973 SPIRV::MemorySemantics::UniformMemory |
4974 SPIRV::MemorySemantics::ImageMemory |
4975 SPIRV::MemorySemantics::WorkgroupMemory,
4977 case Intrinsic::spv_device_memory_barrier:
4978 return selectBarrierInst(
I, SPIRV::Scope::Device,
4979 SPIRV::MemorySemantics::UniformMemory |
4980 SPIRV::MemorySemantics::ImageMemory,
4982 case Intrinsic::spv_device_memory_barrier_with_group_sync:
4983 return selectBarrierInst(
I, SPIRV::Scope::Device,
4984 SPIRV::MemorySemantics::UniformMemory |
4985 SPIRV::MemorySemantics::ImageMemory,
4987 case Intrinsic::spv_group_memory_barrier:
4988 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4989 SPIRV::MemorySemantics::WorkgroupMemory,
4991 case Intrinsic::spv_group_memory_barrier_with_group_sync:
4992 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4993 SPIRV::MemorySemantics::WorkgroupMemory,
4995 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4996 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4997 SPIRV::StorageClass::StorageClass ResSC =
5001 "Generic storage class");
5002 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5010 case Intrinsic::spv_lifetime_start:
5011 case Intrinsic::spv_lifetime_end: {
5012 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5013 : SPIRV::OpLifetimeStop;
5014 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5015 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5024 case Intrinsic::spv_saturate:
5025 return selectSaturate(ResVReg, ResType,
I);
5026 case Intrinsic::spv_nclamp:
5027 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5028 case Intrinsic::spv_uclamp:
5029 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5030 case Intrinsic::spv_sclamp:
5031 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5032 case Intrinsic::spv_subgroup_prefix_bit_count:
5033 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5034 case Intrinsic::spv_wave_active_countbits:
5035 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5036 case Intrinsic::spv_wave_all_equal:
5037 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5038 case Intrinsic::spv_wave_all:
5039 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5040 case Intrinsic::spv_wave_any:
5041 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5042 case Intrinsic::spv_subgroup_ballot:
5043 return selectWaveOpInst(ResVReg, ResType,
I,
5044 SPIRV::OpGroupNonUniformBallot);
5045 case Intrinsic::spv_wave_is_first_lane:
5046 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5047 case Intrinsic::spv_wave_reduce_or:
5048 return selectWaveReduceOp(ResVReg, ResType,
I,
5049 SPIRV::OpGroupNonUniformBitwiseOr);
5050 case Intrinsic::spv_wave_reduce_xor:
5051 return selectWaveReduceOp(ResVReg, ResType,
I,
5052 SPIRV::OpGroupNonUniformBitwiseXor);
5053 case Intrinsic::spv_wave_reduce_and:
5054 return selectWaveReduceOp(ResVReg, ResType,
I,
5055 SPIRV::OpGroupNonUniformBitwiseAnd);
5056 case Intrinsic::spv_wave_reduce_umax:
5057 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5058 case Intrinsic::spv_wave_reduce_max:
5059 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5060 case Intrinsic::spv_wave_reduce_umin:
5061 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5062 case Intrinsic::spv_wave_reduce_min:
5063 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5064 case Intrinsic::spv_wave_reduce_sum:
5065 return selectWaveReduceSum(ResVReg, ResType,
I);
5066 case Intrinsic::spv_wave_product:
5067 return selectWaveReduceProduct(ResVReg, ResType,
I);
5068 case Intrinsic::spv_wave_readlane:
5069 return selectWaveOpInst(ResVReg, ResType,
I,
5070 SPIRV::OpGroupNonUniformShuffle);
5071 case Intrinsic::spv_wave_prefix_sum:
5072 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5073 case Intrinsic::spv_wave_prefix_product:
5074 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5075 case Intrinsic::spv_quad_read_across_x: {
5076 return selectQuadSwap(ResVReg, ResType,
I, 0);
5078 case Intrinsic::spv_quad_read_across_y: {
5079 return selectQuadSwap(ResVReg, ResType,
I, 1);
5081 case Intrinsic::spv_step:
5082 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5083 case Intrinsic::spv_radians:
5084 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5088 case Intrinsic::instrprof_increment:
5089 case Intrinsic::instrprof_increment_step:
5090 case Intrinsic::instrprof_value_profile:
5093 case Intrinsic::spv_value_md:
5095 case Intrinsic::spv_resource_handlefrombinding: {
5096 return selectHandleFromBinding(ResVReg, ResType,
I);
5098 case Intrinsic::spv_resource_counterhandlefrombinding:
5099 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5100 case Intrinsic::spv_resource_updatecounter:
5101 return selectUpdateCounter(ResVReg, ResType,
I);
5102 case Intrinsic::spv_resource_store_typedbuffer: {
5103 return selectImageWriteIntrinsic(
I);
5105 case Intrinsic::spv_resource_load_typedbuffer: {
5106 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5108 case Intrinsic::spv_resource_load_level: {
5109 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5111 case Intrinsic::spv_resource_getdimensions_x:
5112 case Intrinsic::spv_resource_getdimensions_xy:
5113 case Intrinsic::spv_resource_getdimensions_xyz: {
5114 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5116 case Intrinsic::spv_resource_getdimensions_levels_x:
5117 case Intrinsic::spv_resource_getdimensions_levels_xy:
5118 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5119 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5121 case Intrinsic::spv_resource_getdimensions_ms_xy:
5122 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5123 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5125 case Intrinsic::spv_resource_calculate_lod:
5126 case Intrinsic::spv_resource_calculate_lod_unclamped:
5127 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5128 case Intrinsic::spv_resource_sample:
5129 case Intrinsic::spv_resource_sample_clamp:
5130 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5131 case Intrinsic::spv_resource_samplebias:
5132 case Intrinsic::spv_resource_samplebias_clamp:
5133 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5134 case Intrinsic::spv_resource_samplegrad:
5135 case Intrinsic::spv_resource_samplegrad_clamp:
5136 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5137 case Intrinsic::spv_resource_samplelevel:
5138 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5139 case Intrinsic::spv_resource_samplecmp:
5140 case Intrinsic::spv_resource_samplecmp_clamp:
5141 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5142 case Intrinsic::spv_resource_samplecmplevelzero:
5143 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5144 case Intrinsic::spv_resource_gather:
5145 case Intrinsic::spv_resource_gather_cmp:
5146 return selectGatherIntrinsic(ResVReg, ResType,
I);
5147 case Intrinsic::spv_resource_getbasepointer:
5148 case Intrinsic::spv_resource_getpointer: {
5149 return selectResourceGetPointer(ResVReg, ResType,
I);
5151 case Intrinsic::spv_pushconstant_getpointer: {
5152 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5154 case Intrinsic::spv_discard: {
5155 return selectDiscard(ResVReg, ResType,
I);
5157 case Intrinsic::spv_resource_nonuniformindex: {
5158 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5160 case Intrinsic::spv_unpackhalf2x16: {
5161 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5163 case Intrinsic::spv_packhalf2x16: {
5164 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5166 case Intrinsic::spv_ddx:
5167 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5168 case Intrinsic::spv_ddy:
5169 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5170 case Intrinsic::spv_ddx_coarse:
5171 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5172 case Intrinsic::spv_ddy_coarse:
5173 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5174 case Intrinsic::spv_ddx_fine:
5175 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5176 case Intrinsic::spv_ddy_fine:
5177 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5178 case Intrinsic::spv_fwidth:
5179 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5180 case Intrinsic::spv_masked_gather:
5181 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5182 return selectMaskedGather(ResVReg, ResType,
I);
5183 return diagnoseUnsupported(
5184 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5185 case Intrinsic::spv_masked_scatter:
5186 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5187 return selectMaskedScatter(
I);
5188 return diagnoseUnsupported(
5189 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5190 case Intrinsic::returnaddress:
5191 case Intrinsic::frameaddress: {
5193 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5200 std::string DiagMsg;
5201 raw_string_ostream OS(DiagMsg);
5203 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
5210bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5211 SPIRVTypeInst ResType,
5212 MachineInstr &
I)
const {
5215 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5222bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5223 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5225 assert(Intr.getIntrinsicID() ==
5226 Intrinsic::spv_resource_counterhandlefrombinding);
5229 Register MainHandleReg = Intr.getOperand(2).getReg();
5231 assert(MainHandleDef->getIntrinsicID() ==
5232 Intrinsic::spv_resource_handlefrombinding);
5236 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5237 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5238 std::string CounterName =
5243 MachineIRBuilder MIRBuilder(
I);
5245 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5247 ArraySize, IndexReg, CounterName, MIRBuilder);
5249 return BuildCOPY(ResVReg, CounterVarReg,
I);
5252bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5253 SPIRVTypeInst ResType,
5254 MachineInstr &
I)
const {
5256 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5258 Register CounterHandleReg = Intr.getOperand(2).getReg();
5259 Register IncrReg = Intr.getOperand(3).getReg();
5266 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5267 assert(CounterVarPointeeType &&
5268 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5269 "Counter variable must be a struct");
5271 SPIRV::StorageClass::StorageBuffer &&
5272 "Counter variable must be in the storage buffer storage class");
5274 "Counter variable must have exactly 1 member in the struct");
5275 const SPIRVTypeInst MemberType =
5278 "Counter variable struct must have a single i32 member");
5282 MachineIRBuilder MIRBuilder(
I);
5284 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5287 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5293 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5296 .
addUse(CounterHandleReg)
5303 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5306 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5309 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5318 return BuildCOPY(ResVReg, AtomicRes,
I);
5326 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5334bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5335 SPIRVTypeInst ResType,
5336 MachineInstr &
I)
const {
5344 Register ImageReg =
I.getOperand(2).getReg();
5352 Register IdxReg =
I.getOperand(3).getReg();
5354 MachineInstr &Pos =
I;
5356 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5360bool SPIRVInstructionSelector::generateSampleImage(
5363 DebugLoc Loc, MachineInstr &Pos)
const {
5374 if (!loadHandleBeforePosition(NewSamplerReg,
5380 MachineIRBuilder MIRBuilder(Pos);
5393 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5394 ImOps.Lod.has_value();
5395 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5396 : SPIRV::OpImageSampleImplicitLod;
5398 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5399 : SPIRV::OpImageSampleDrefImplicitLod;
5408 MIB.
addUse(*ImOps.Compare);
5410 uint32_t ImageOperands = 0;
5412 ImageOperands |= SPIRV::ImageOperand::Bias;
5414 ImageOperands |= SPIRV::ImageOperand::Lod;
5415 if (ImOps.GradX && ImOps.GradY)
5416 ImageOperands |= SPIRV::ImageOperand::Grad;
5417 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5419 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5422 "Non-constant offsets are not supported in sample instructions.");
5426 ImageOperands |= SPIRV::ImageOperand::MinLod;
5428 if (ImageOperands != 0) {
5429 MIB.
addImm(ImageOperands);
5430 if (ImageOperands & SPIRV::ImageOperand::Bias)
5432 if (ImageOperands & SPIRV::ImageOperand::Lod)
5434 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5435 MIB.
addUse(*ImOps.GradX);
5436 MIB.
addUse(*ImOps.GradY);
5439 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5440 MIB.
addUse(*ImOps.Offset);
5441 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5442 MIB.
addUse(*ImOps.MinLod);
5449bool SPIRVInstructionSelector::selectImageQuerySize(
5451 std::optional<Register> LodReg)
const {
5453 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5456 "ImageReg is not an image type.");
5458 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5460 unsigned NumComponents = 0;
5462 case SPIRV::Dim::DIM_1D:
5463 case SPIRV::Dim::DIM_Buffer:
5464 NumComponents =
IsArray ? 2 : 1;
5466 case SPIRV::Dim::DIM_2D:
5467 case SPIRV::Dim::DIM_Cube:
5468 case SPIRV::Dim::DIM_Rect:
5469 NumComponents =
IsArray ? 3 : 2;
5471 case SPIRV::Dim::DIM_3D:
5475 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5480 SPIRVTypeInst ResType =
5485 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5495bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5496 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5497 Register ImageReg =
I.getOperand(2).getReg();
5504 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5507bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5508 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5509 Register ImageReg =
I.getOperand(2).getReg();
5518 Register LodReg =
I.getOperand(3).getReg();
5521 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5523 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5530 TII.get(SPIRV::OpImageQueryLevels))
5537 TII.get(SPIRV::OpCompositeConstruct))
5547bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5548 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5549 Register ImageReg =
I.getOperand(2).getReg();
5560 "OpImageQuerySamples requires a multisampled image");
5562 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5570 TII.get(SPIRV::OpImageQuerySamples))
5577 TII.get(SPIRV::OpCompositeConstruct))
5587bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5588 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5589 Register ImageReg =
I.getOperand(2).getReg();
5590 Register SamplerReg =
I.getOperand(3).getReg();
5591 Register CoordinateReg =
I.getOperand(4).getReg();
5607 if (!loadHandleBeforePosition(
5612 MachineIRBuilder MIRBuilder(
I);
5618 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5628 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
5635 unsigned ExtractedIndex =
5637 Intrinsic::spv_resource_calculate_lod_unclamped
5641 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5642 TII.get(SPIRV::OpCompositeExtract))
5652bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
5653 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5654 Register ImageReg =
I.getOperand(2).getReg();
5655 Register SamplerReg =
I.getOperand(3).getReg();
5656 Register CoordinateReg =
I.getOperand(4).getReg();
5657 ImageOperands ImOps;
5658 if (
I.getNumOperands() > 5)
5659 ImOps.Offset =
I.getOperand(5).getReg();
5660 if (
I.getNumOperands() > 6)
5661 ImOps.MinLod =
I.getOperand(6).getReg();
5662 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5663 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5666bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5667 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5668 Register ImageReg =
I.getOperand(2).getReg();
5669 Register SamplerReg =
I.getOperand(3).getReg();
5670 Register CoordinateReg =
I.getOperand(4).getReg();
5671 ImageOperands ImOps;
5672 ImOps.Bias =
I.getOperand(5).getReg();
5673 if (
I.getNumOperands() > 6)
5674 ImOps.Offset =
I.getOperand(6).getReg();
5675 if (
I.getNumOperands() > 7)
5676 ImOps.MinLod =
I.getOperand(7).getReg();
5677 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5678 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5681bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
5682 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5683 Register ImageReg =
I.getOperand(2).getReg();
5684 Register SamplerReg =
I.getOperand(3).getReg();
5685 Register CoordinateReg =
I.getOperand(4).getReg();
5686 ImageOperands ImOps;
5687 ImOps.GradX =
I.getOperand(5).getReg();
5688 ImOps.GradY =
I.getOperand(6).getReg();
5689 if (
I.getNumOperands() > 7)
5690 ImOps.Offset =
I.getOperand(7).getReg();
5691 if (
I.getNumOperands() > 8)
5692 ImOps.MinLod =
I.getOperand(8).getReg();
5693 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5694 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5697bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
5698 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5699 Register ImageReg =
I.getOperand(2).getReg();
5700 Register SamplerReg =
I.getOperand(3).getReg();
5701 Register CoordinateReg =
I.getOperand(4).getReg();
5702 ImageOperands ImOps;
5703 ImOps.Lod =
I.getOperand(5).getReg();
5704 if (
I.getNumOperands() > 6)
5705 ImOps.Offset =
I.getOperand(6).getReg();
5706 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5707 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5710bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
5711 SPIRVTypeInst ResType,
5712 MachineInstr &
I)
const {
5713 Register ImageReg =
I.getOperand(2).getReg();
5714 Register SamplerReg =
I.getOperand(3).getReg();
5715 Register CoordinateReg =
I.getOperand(4).getReg();
5716 ImageOperands ImOps;
5717 ImOps.Compare =
I.getOperand(5).getReg();
5718 if (
I.getNumOperands() > 6)
5719 ImOps.Offset =
I.getOperand(6).getReg();
5720 if (
I.getNumOperands() > 7)
5721 ImOps.MinLod =
I.getOperand(7).getReg();
5722 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5723 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5726bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
5727 SPIRVTypeInst ResType,
5728 MachineInstr &
I)
const {
5729 Register ImageReg =
I.getOperand(2).getReg();
5730 Register CoordinateReg =
I.getOperand(3).getReg();
5731 Register LodReg =
I.getOperand(4).getReg();
5733 ImageOperands ImOps;
5735 if (
I.getNumOperands() > 5)
5736 ImOps.Offset =
I.getOperand(5).getReg();
5748 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
5749 I.getDebugLoc(),
I, &ImOps);
5752bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
5753 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5754 Register ImageReg =
I.getOperand(2).getReg();
5755 Register SamplerReg =
I.getOperand(3).getReg();
5756 Register CoordinateReg =
I.getOperand(4).getReg();
5757 ImageOperands ImOps;
5758 ImOps.Compare =
I.getOperand(5).getReg();
5759 if (
I.getNumOperands() > 6)
5760 ImOps.Offset =
I.getOperand(6).getReg();
5763 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5764 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5767bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
5768 SPIRVTypeInst ResType,
5769 MachineInstr &
I)
const {
5770 Register ImageReg =
I.getOperand(2).getReg();
5771 Register SamplerReg =
I.getOperand(3).getReg();
5772 Register CoordinateReg =
I.getOperand(4).getReg();
5775 "ImageReg is not an image type.");
5780 ComponentOrCompareReg =
I.getOperand(5).getReg();
5781 OffsetReg =
I.getOperand(6).getReg();
5784 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
5788 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5789 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
5790 Dim != SPIRV::Dim::DIM_Rect) {
5792 "Gather operations are only supported for 2D, Cube, and Rect images.");
5799 if (!loadHandleBeforePosition(
5804 MachineIRBuilder MIRBuilder(
I);
5805 SPIRVTypeInst SampledImageType =
5810 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5818 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
5820 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
5822 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5827 .
addUse(ComponentOrCompareReg);
5829 uint32_t ImageOperands = 0;
5830 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
5831 if (Dim == SPIRV::Dim::DIM_Cube) {
5833 "Gather operations with offset are not supported for Cube images.");
5837 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5839 ImageOperands |= SPIRV::ImageOperand::Offset;
5843 if (ImageOperands != 0) {
5844 MIB.
addImm(ImageOperands);
5846 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5854bool SPIRVInstructionSelector::generateImageReadOrFetch(
5857 const ImageOperands *ImOps)
const {
5860 "ImageReg is not an image type.");
5862 bool IsSignedInteger =
5867 bool IsFetch = (SampledOp.getImm() == 1);
5869 auto AddOperands = [&](MachineInstrBuilder &MIB) {
5870 uint32_t ImageOperandsMask = 0;
5871 if (IsSignedInteger)
5872 ImageOperandsMask |= 0x1000;
5874 if (IsFetch && ImOps) {
5876 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
5877 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
5879 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
5881 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5885 if (ImageOperandsMask != 0) {
5886 MIB.
addImm(ImageOperandsMask);
5887 if (IsFetch && ImOps) {
5890 if (ImOps->Offset &&
5891 (ImageOperandsMask &
5892 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5893 MIB.
addUse(*ImOps->Offset);
5899 if (ResultSize == 4) {
5902 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5909 BMI.constrainAllUses(
TII,
TRI, RBI);
5913 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5917 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5923 BMI.constrainAllUses(
TII,
TRI, RBI);
5925 if (ResultSize == 1) {
5934 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
5937bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
5938 SPIRVTypeInst ResType,
5939 MachineInstr &
I)
const {
5940 Register ResourcePtr =
I.getOperand(2).getReg();
5942 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
5951 MachineIRBuilder MIRBuilder(
I);
5956 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5962 if (
I.getNumExplicitOperands() > 3) {
5963 Register IndexReg =
I.getOperand(3).getReg();
5970bool SPIRVInstructionSelector::selectPushConstantGetPointer(
5971 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5976bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
5977 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5978 Register ObjReg =
I.getOperand(2).getReg();
5979 if (!BuildCOPY(ResVReg, ObjReg,
I))
5989 decorateUsesAsNonUniform(ResVReg);
5993void SPIRVInstructionSelector::decorateUsesAsNonUniform(
5996 while (WorkList.
size() > 0) {
6000 bool IsDecorated =
false;
6002 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6003 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6009 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6011 if (ResultReg == CurrentReg)
6019 SPIRV::Decoration::NonUniformEXT, {});
6024bool SPIRVInstructionSelector::extractSubvector(
6026 MachineInstr &InsertionPoint)
const {
6028 [[maybe_unused]] uint64_t InputSize =
6031 assert(InputSize > 1 &&
"The input must be a vector.");
6032 assert(ResultSize > 1 &&
"The result must be a vector.");
6033 assert(ResultSize < InputSize &&
6034 "Cannot extract more element than there are in the input.");
6037 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
6038 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6041 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6050 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6052 TII.get(SPIRV::OpCompositeConstruct))
6056 for (
Register ComponentReg : ComponentRegisters)
6057 MIB.
addUse(ComponentReg);
6062bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6063 MachineInstr &
I)
const {
6070 Register ImageReg =
I.getOperand(1).getReg();
6078 Register CoordinateReg =
I.getOperand(2).getReg();
6079 Register DataReg =
I.getOperand(3).getReg();
6082 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6090Register SPIRVInstructionSelector::buildPointerToResource(
6091 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6092 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6093 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6095 if (ArraySize == 1) {
6096 SPIRVTypeInst PtrType =
6099 "SpirvResType did not have an explicit layout.");
6104 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6105 SPIRVTypeInst VarPointerType =
6108 VarPointerType, Set,
Binding, Name, MIRBuilder);
6110 SPIRVTypeInst ResPointerType =
6123bool SPIRVInstructionSelector::selectFirstBitSet16(
6124 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6125 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6127 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6131 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6134bool SPIRVInstructionSelector::selectFirstBitSet32(
6136 unsigned BitSetOpcode)
const {
6137 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6140 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6147bool SPIRVInstructionSelector::selectFirstBitSet64(
6149 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6162 if (ComponentCount > 2) {
6163 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6165 unsigned Opcode) ->
bool {
6166 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6170 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6174 MachineIRBuilder MIRBuilder(
I);
6176 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6180 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6186 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6193 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6196 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6197 SPIRV::OpVectorExtractDynamic))
6199 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6200 SPIRV::OpVectorExtractDynamic))
6204 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6205 TII.get(SPIRV::OpVectorShuffle))
6213 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6219 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6220 TII.get(SPIRV::OpVectorShuffle))
6228 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6248 SelectOp = SPIRV::OpSelectSISCond;
6249 AddOp = SPIRV::OpIAddS;
6257 SelectOp = SPIRV::OpSelectVIVCond;
6258 AddOp = SPIRV::OpIAddV;
6264 Register RegSecondaryOffset = Reg0;
6268 if (SwapPrimarySide) {
6269 PrimaryReg = LowReg;
6270 SecondaryReg = HighReg;
6271 RegPrimaryOffset = Reg0;
6272 RegSecondaryOffset = Reg32;
6277 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6278 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6283 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6284 SPIRV::OpINotEqual))
6291 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6292 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6297 if (SwapPrimarySide) {
6299 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6300 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6311 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6312 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6317 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6318 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6321 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6325bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6326 SPIRVTypeInst ResType,
6328 bool IsSigned)
const {
6330 Register OpReg =
I.getOperand(2).getReg();
6333 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6334 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6338 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6340 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6342 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6346 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6350bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6351 SPIRVTypeInst ResType,
6352 MachineInstr &
I)
const {
6354 Register OpReg =
I.getOperand(2).getReg();
6359 unsigned ExtendOpcode = SPIRV::OpUConvert;
6360 unsigned BitSetOpcode = GL::FindILsb;
6364 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6366 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6368 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6375bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6376 SPIRVTypeInst ResType,
6377 MachineInstr &
I)
const {
6381 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6384 .
addUse(
I.getOperand(2).getReg())
6387 unsigned Alignment =
I.getOperand(3).getImm();
6401 while (!Worklist.
empty()) {
6403 switch (
T->getOpcode()) {
6404 case SPIRV::OpTypeInt:
6405 case SPIRV::OpTypeFloat:
6406 case SPIRV::OpTypePointer:
6408 case SPIRV::OpTypeVector:
6409 case SPIRV::OpTypeMatrix:
6410 case SPIRV::OpTypeArray: {
6411 Register OperandReg =
T->getOperand(1).getReg();
6415 case SPIRV::OpTypeStruct:
6416 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6417 Register OperandReg =
T->getOperand(Idx).getReg();
6429bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6430 assert(
I.getNumExplicitOperands() == 2);
6432 Register MsgReg =
I.getOperand(1).getReg();
6434 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6437 return diagnoseUnsupported(
6439 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6440 "scalar, pointer, vector, matrix, or aggregate of such types)");
6443 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6450bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6459 uint32_t MsgVal = ~0
u;
6460 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6461 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6464 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6467 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6474bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6475 SPIRVTypeInst ResType,
6476 MachineInstr &
I)
const {
6480 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6483 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6486 unsigned Alignment =
I.getOperand(2).getImm();
6493bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6498 const MachineInstr *PrevI =
I.getPrevNode();
6500 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6504 .
addMBB(
I.getOperand(0).getMBB())
6509 .
addMBB(
I.getOperand(0).getMBB())
6514bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6525 const MachineInstr *NextI =
I.getNextNode();
6527 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6533 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6535 .
addUse(
I.getOperand(0).getReg())
6536 .
addMBB(
I.getOperand(1).getMBB())
6542bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6543 MachineInstr &
I)
const {
6545 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6547 const unsigned NumOps =
I.getNumOperands();
6548 for (
unsigned i = 1; i <
NumOps; i += 2) {
6549 MIB.
addUse(
I.getOperand(i + 0).getReg());
6550 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6556bool SPIRVInstructionSelector::selectGlobalValue(
6557 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6559 MachineIRBuilder MIRBuilder(
I);
6560 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6563 std::string GlobalIdent;
6565 unsigned &
ID = UnnamedGlobalIDs[GV];
6567 ID = UnnamedGlobalIDs.
size();
6568 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6594 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6601 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6606 MachineInstrBuilder MIB1 =
6607 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6610 MachineInstrBuilder MIB2 =
6612 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
6616 GR.
add(ConstVal, MIB2);
6624 MachineInstrBuilder MIB3 =
6625 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
6628 GR.
add(ConstVal, MIB3);
6632 assert(NewReg != ResVReg);
6633 return BuildCOPY(ResVReg, NewReg,
I);
6643 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
6646 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
6652 SPIRVTypeInst ResType =
6656 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6661 if (
GlobalVar->isExternallyInitialized() &&
6662 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6663 constexpr unsigned ReadWriteINTEL = 3u;
6666 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6672bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6673 SPIRVTypeInst ResType,
6674 MachineInstr &
I)
const {
6676 return selectExtInst(ResVReg, ResType,
I, CL::log10);
6684 MachineIRBuilder MIRBuilder(
I);
6689 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6692 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6694 .
add(
I.getOperand(1))
6699 ResType->
getOpcode() == SPIRV::OpTypeFloat);
6709 APFloat::rmNearestTiesToEven, &LosesInfo);
6713 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
6714 ? SPIRV::OpVectorTimesScalar
6725bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
6726 SPIRVTypeInst ResType,
6727 MachineInstr &
I)
const {
6730 return selectExtInst(ResVReg, ResType,
I, CL::pown);
6736 Register ExpReg =
I.getOperand(2).getReg();
6738 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
6739 SPIRV::OpConvertSToF))
6741 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
6748bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
6749 SPIRVTypeInst ResType,
6750 MachineInstr &
I)
const {
6766 MachineIRBuilder MIRBuilder(
I);
6769 ResType, MIRBuilder, SPIRV::StorageClass::Function);
6782 MachineBasicBlock &EntryBB =
I.getMF()->
front();
6784 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
6787 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
6793 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6796 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
6799 .
add(
I.getOperand(
I.getNumExplicitDefs()))
6803 Register IntegralPartReg =
I.getOperand(1).getReg();
6804 if (IntegralPartReg.
isValid()) {
6806 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6817 assert(
false &&
"GLSL::Modf is deprecated.");
6828bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
6829 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6830 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6831 MachineIRBuilder MIRBuilder(
I);
6832 const SPIRVTypeInst Vec3Ty =
6835 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
6847 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6851 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6857 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6864 assert(
I.getOperand(2).isReg());
6865 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
6869 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6880bool SPIRVInstructionSelector::loadBuiltinInputID(
6881 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6882 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6883 MachineIRBuilder MIRBuilder(
I);
6885 ResType, MIRBuilder, SPIRV::StorageClass::Input);
6900 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6904 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6913SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
6914 MachineInstr &
I)
const {
6915 MachineIRBuilder MIRBuilder(
I);
6916 if (
Type->getOpcode() != SPIRV::OpTypeVector)
6926bool SPIRVInstructionSelector::loadHandleBeforePosition(
6927 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
6928 MachineInstr &Pos)
const {
6931 Intrinsic::spv_resource_handlefrombinding);
6939 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
6940 MachineIRBuilder MIRBuilder(HandleDef);
6941 SPIRVTypeInst VarType = ResType;
6942 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
6944 if (IsStructuredBuffer) {
6949 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
6951 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
6954 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
6955 ArraySize, IndexReg, Name, MIRBuilder);
6959 uint32_t LoadOpcode =
6960 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
6970void SPIRVInstructionSelector::errorIfInstrOutsideShader(
6971 MachineInstr &
I)
const {
6973 std::string DiagMsg;
6974 raw_string_ostream OS(DiagMsg);
6975 I.print(OS,
true,
false,
false,
false);
6976 DiagMsg +=
" is only supported in shaders.\n";
6982InstructionSelector *
6986 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
MachineInstrBuilder & UseMI
#define GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
This file declares a class to represent arbitrary precision floating point values and provide a varie...
static bool selectUnmergeValues(MachineInstrBuilder &MIB, const ARMBaseInstrInfo &TII, MachineRegisterInfo &MRI, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
static uint8_t SwapBits(uint8_t Val)
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
DXIL Resource Implicit Binding
Declares convenience wrapper classes for interpreting MachineInstr instances as specific generic oper...
const HexagonInstrInfo * TII
LLVMTypeRef LLVMIntType(unsigned NumBits)
const size_t AbstractManglingParser< Derived, Alloc >::NumOps
Loop::LoopBounds::Direction Direction
Register const TargetRegisterInfo * TRI
Promote Memory to Register
MachineInstr unsigned OpIdx
uint64_t IntrinsicInst * II
static StringRef getName(Value *V)
static unsigned getFCmpOpcode(CmpInst::Predicate Pred, unsigned Size)
static bool isConcreteSPIRVType(SPIRVTypeInst Ty, const SPIRVGlobalRegistry &GR)
static APFloat getOneFP(const Type *LLVMFloatTy)
static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC)
static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg)
static bool mayApplyGenericSelection(unsigned Opcode)
static APFloat getZeroFP(const Type *LLVMFloatTy)
std::vector< std::pair< SPIRV::InstructionSet::InstructionSet, uint32_t > > ExtInstList
static bool intrinsicHasSideEffects(Intrinsic::ID ID)
static unsigned getBoolCmpOpcode(unsigned PredNum)
static unsigned getICmpOpcode(unsigned PredNum)
static bool isOpcodeWithNoSideEffects(unsigned Opcode)
static void addMemoryOperands(MachineMemOperand *MemOp, MachineInstrBuilder &MIB, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry &GR)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef)
static unsigned getPtrCmpOpcode(unsigned Pred)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
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 & 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.
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.
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.
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 emitGenericError(const Twine &ErrMsg) const
LLVM_ABI const MachineFunction * getMF() const
Return the function that contains the basic block that this instruction belongs to.
const DebugLoc & getDebugLoc() const
Returns the debug location id of this MachineInstr.
const MachineOperand & getOperand(unsigned i) const
A description of a memory reference used in the backend.
@ MOVolatile
The memory access is volatile.
@ MONonTemporal
The memory access is non-temporal.
bool isReg() const
isReg - Tests if this is a MO_Register operand.
MachineBasicBlock * getMBB() const
Register getReg() const
getReg - Returns the register number.
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
defusechain_instr_iterator< true, false, false, true > use_instr_iterator
use_instr_iterator/use_instr_begin/use_instr_end - Walk all uses of the specified register,...
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
LLVM_ABI 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
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.
static use_instr_iterator use_instr_end()
iterator_range< use_instr_nodbg_iterator > use_nodbg_instructions(Register Reg) const
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
const MachineFunction & getMF() const
LLVM_ABI void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
LLVM_ABI Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
const TargetRegisterClass * getRegClassOrNull(Register Reg) const
Return the register class of Reg, or null if Reg has not been assigned a register class yet.
iterator_range< use_instr_iterator > use_instructions(Register Reg) const
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
LLVM_ABI void replaceRegWith(Register FromReg, Register ToReg)
replaceRegWith - Replace all instances of FromReg with ToReg in the machine function.
Analysis providing profile information.
Holds all the information related to register banks.
Wrapper class representing virtual and physical registers.
constexpr bool isValid() const
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
bool isScalarOrVectorSigned(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateOpTypeSampledImage(SPIRVTypeInst ImageType, MachineIRBuilder &MIRBuilder)
void assignSPIRVTypeToVReg(SPIRVTypeInst Type, Register VReg, const MachineFunction &MF)
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
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC)
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
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 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.
reference emplace_back(ArgTypes &&... Args)
void reserve(size_type N)
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
constexpr size_t size() const
Get the string size.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
@ HalfTyID
16-bit floating point type
@ FloatTyID
32-bit floating point type
@ 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.
self_iterator getIterator()
NodeTy * getNextNode()
Get the next node, or nullptr for the list tail.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
constexpr char IsConst[]
Key for Kernel::Arg::Metadata::mIsConst.
constexpr std::underlying_type_t< E > Mask()
Get a bitmask with 1s in all places up to the high-order bit of E's largest value.
unsigned ID
LLVM IR allows to use arbitrary numbers as calling convention identifiers.
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.
This is an optimization pass for GlobalISel generic memory operations.
void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)
@ Low
Lower the current thread's priority such that it does not affect foreground tasks significantly.
FunctionAddr VTableAddr Value
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
MachineBasicBlock::iterator getOpVariableMBBIt(MachineFunction &MF)
int64_t getIConstValSext(Register ConstReg, const MachineRegisterInfo *MRI)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
FunctionAddr VTableAddr uintptr_t uintptr_t Int32Ty
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.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
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 any_of(R &&range, UnaryPredicate P)
Provide wrappers to std::any_of which take ranges instead of having to pass begin/end explicitly.
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)
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.
SPIRV::Scope::Scope getMemScope(LLVMContext &Ctx, SyncScope::ID Id)
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)
void addStringImm(const StringRef &Str, MCInst &Inst)
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...