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();
3975 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3979SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3980 SPIRVTypeInst ResType)
const {
3982 SPIRVTypeInst SpvI32Ty =
3985 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3992 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3995 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3998 .
addImm(APInt(32, Val).getZExtValue());
4000 GR.
add(ConstInt,
MI);
4007Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4008 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4010 SPIRVTypeInst SpvI32Ty =
4012 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4017 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4018 MachineInstr *
MI =
nullptr;
4022 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4026 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4027 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4033 GR.
add(ConstInt,
MI);
4038bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4039 SPIRVTypeInst ResType,
4040 MachineInstr &
I)
const {
4042 return selectCmp(ResVReg, ResType, CmpOp,
I);
4045bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4046 SPIRVTypeInst ResType,
4047 MachineInstr &
I)
const {
4049 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4056 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4057 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4060 MachineIRBuilder MIRBuilder(
I);
4067 APFloat ConstVal(3.3219280948873623);
4071 APFloat::rmNearestTiesToEven, &LosesInfo);
4075 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4076 ? SPIRV::OpVectorTimesScalar
4079 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4080 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4082 if (!selectExtInst(ResVReg, ResType,
I,
4083 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4093Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4094 MachineInstr &
I)
const {
4097 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4102bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4108 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4116 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4119 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4120 Def->getOpcode() == SPIRV::OpConstantI)
4133 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4134 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4136 Intrinsic::spv_const_composite)) {
4137 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4138 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4139 if (!IsZero(
Def->getOperand(i).getReg()))
4148Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4149 MachineInstr &
I)
const {
4153 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4158Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4159 MachineInstr &
I)
const {
4163 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4169 SPIRVTypeInst ResType,
4170 MachineInstr &
I)
const {
4174 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4179bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4180 SPIRVTypeInst ResType,
4181 MachineInstr &
I)
const {
4182 Register SelectFirstArg =
I.getOperand(2).getReg();
4183 Register SelectSecondArg =
I.getOperand(3).getReg();
4192 SPIRV::OpTypeVector;
4199 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4200 }
else if (IsPtrTy) {
4201 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4203 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4206 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4207 "boolean condition");
4209 Opcode = SPIRV::OpSelectSFSCond;
4210 }
else if (IsPtrTy) {
4211 Opcode = SPIRV::OpSelectSPSCond;
4213 Opcode = SPIRV::OpSelectSISCond;
4216 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4219 .
addUse(
I.getOperand(1).getReg())
4228bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4229 SPIRVTypeInst ResType,
4231 MachineInstr &InsertAt,
4232 bool IsSigned)
const {
4234 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4235 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4236 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4238 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4250bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4251 SPIRVTypeInst ResType,
4252 MachineInstr &
I,
bool IsSigned,
4253 unsigned Opcode)
const {
4254 Register SrcReg =
I.getOperand(1).getReg();
4260 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4265 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
4267 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4270bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4271 SPIRVTypeInst ResType, MachineInstr &
I,
4272 bool IsSigned)
const {
4273 Register SrcReg =
I.getOperand(1).getReg();
4275 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4279 if (ResType == SrcType)
4280 return BuildCOPY(ResVReg, SrcReg,
I);
4282 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4283 return selectUnOp(ResVReg, ResType,
I, Opcode);
4286bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4287 SPIRVTypeInst ResType,
4289 bool IsSigned)
const {
4290 MachineIRBuilder MIRBuilder(
I);
4291 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4303 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4306 .
addUse(
I.getOperand(1).getReg())
4307 .
addUse(
I.getOperand(2).getReg())
4312 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4315 .
addUse(
I.getOperand(1).getReg())
4316 .
addUse(
I.getOperand(2).getReg())
4324 unsigned SelectOpcode =
4325 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4330 .
addUse(buildOnesVal(
true, ResType,
I))
4331 .
addUse(buildZerosVal(ResType,
I))
4338 .
addUse(buildOnesVal(
false, ResType,
I))
4343bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4346 SPIRVTypeInst IntTy,
4347 SPIRVTypeInst BoolTy)
const {
4350 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4351 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4353 Register One = buildOnesVal(
false, IntTy,
I);
4361 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4370bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4371 SPIRVTypeInst ResType,
4372 MachineInstr &
I)
const {
4373 Register IntReg =
I.getOperand(1).getReg();
4376 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4377 if (ArgType == ResType)
4378 return BuildCOPY(ResVReg, IntReg,
I);
4380 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4381 return selectUnOp(ResVReg, ResType,
I, Opcode);
4384bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4385 SPIRVTypeInst ResType,
4386 MachineInstr &
I)
const {
4387 unsigned Opcode =
I.getOpcode();
4388 unsigned TpOpcode = ResType->
getOpcode();
4390 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4391 assert(Opcode == TargetOpcode::G_CONSTANT &&
4392 I.getOperand(1).getCImm()->isZero());
4393 MachineBasicBlock &DepMBB =
I.getMF()->front();
4396 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4403 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4406bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4407 SPIRVTypeInst ResType,
4408 MachineInstr &
I)
const {
4409 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4416bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4417 SPIRVTypeInst ResType,
4418 MachineInstr &
I)
const {
4420 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4424 .
addUse(
I.getOperand(3).getReg())
4426 .
addUse(
I.getOperand(2).getReg());
4427 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4433bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4434 SPIRVTypeInst ResType,
4435 MachineInstr &
I)
const {
4436 Type *MaybeResTy =
nullptr;
4441 "Expected aggregate type for extractv instruction");
4443 SPIRV::AccessQualifier::ReadWrite,
false);
4447 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4450 .
addUse(
I.getOperand(2).getReg());
4451 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4457bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4458 SPIRVTypeInst ResType,
4459 MachineInstr &
I)
const {
4460 if (
getImm(
I.getOperand(4), MRI))
4461 return selectInsertVal(ResVReg, ResType,
I);
4463 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4466 .
addUse(
I.getOperand(2).getReg())
4467 .
addUse(
I.getOperand(3).getReg())
4468 .
addUse(
I.getOperand(4).getReg())
4473bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4474 SPIRVTypeInst ResType,
4475 MachineInstr &
I)
const {
4476 if (
getImm(
I.getOperand(3), MRI))
4477 return selectExtractVal(ResVReg, ResType,
I);
4479 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4482 .
addUse(
I.getOperand(2).getReg())
4483 .
addUse(
I.getOperand(3).getReg())
4488bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4489 SPIRVTypeInst ResType,
4490 MachineInstr &
I)
const {
4491 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4497 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4498 : SPIRV::OpAccessChain)
4499 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4500 :
SPIRV::OpPtrAccessChain);
4502 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4506 .
addUse(
I.getOperand(3).getReg());
4508 (Opcode == SPIRV::OpPtrAccessChain ||
4509 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4510 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4511 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4514 const unsigned StartingIndex =
4515 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4518 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4519 Res.addUse(
I.getOperand(i).getReg());
4520 Res.constrainAllUses(
TII,
TRI, RBI);
4525bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4527 unsigned Lim =
I.getNumExplicitOperands();
4528 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4529 Register OpReg =
I.getOperand(i).getReg();
4530 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4532 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4533 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4534 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4541 MachineFunction *MF =
I.getMF();
4553 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4554 TII.get(SPIRV::OpSpecConstantOp))
4557 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4559 GR.
add(OpDefine, MIB);
4565bool SPIRVInstructionSelector::selectDerivativeInst(
4566 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4567 const unsigned DPdOpCode)
const {
4570 errorIfInstrOutsideShader(
I);
4575 Register SrcReg =
I.getOperand(2).getReg();
4580 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4583 .
addUse(
I.getOperand(2).getReg());
4585 MachineIRBuilder MIRBuilder(
I);
4588 if (componentCount != 1)
4592 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4596 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4601 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4606 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4614bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4615 SPIRVTypeInst ResType,
4616 MachineInstr &
I)
const {
4620 case Intrinsic::spv_load:
4621 return selectLoad(ResVReg, ResType,
I);
4622 case Intrinsic::spv_atomic_load:
4623 return selectAtomicLoad(ResVReg, ResType,
I);
4624 case Intrinsic::spv_store:
4625 return selectStore(
I);
4626 case Intrinsic::spv_atomic_store:
4627 return selectAtomicStore(
I);
4628 case Intrinsic::spv_extractv:
4629 return selectExtractVal(ResVReg, ResType,
I);
4630 case Intrinsic::spv_insertv:
4631 return selectInsertVal(ResVReg, ResType,
I);
4632 case Intrinsic::spv_extractelt:
4633 return selectExtractElt(ResVReg, ResType,
I);
4634 case Intrinsic::spv_insertelt:
4635 return selectInsertElt(ResVReg, ResType,
I);
4636 case Intrinsic::spv_gep:
4637 return selectGEP(ResVReg, ResType,
I);
4638 case Intrinsic::spv_bitcast: {
4639 Register OpReg =
I.getOperand(2).getReg();
4640 SPIRVTypeInst OpType =
4644 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4646 case Intrinsic::spv_unref_global:
4647 case Intrinsic::spv_init_global: {
4648 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4653 Register GVarVReg =
MI->getOperand(0).getReg();
4654 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4659 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4661 MI->eraseFromParent();
4665 case Intrinsic::spv_undef: {
4666 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4672 case Intrinsic::spv_named_boolean_spec_constant: {
4673 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4674 : SPIRV::OpSpecConstantFalse;
4676 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4677 .
addDef(
I.getOperand(0).getReg())
4680 unsigned SpecId =
I.getOperand(2).getImm();
4682 SPIRV::Decoration::SpecId, {SpecId});
4686 case Intrinsic::spv_const_composite: {
4688 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4694 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4696 std::function<bool(
Register)> HasSpecConstOperand =
4706 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
4707 J < Def->getNumExplicitOperands(); ++J) {
4708 if (
Def->getOperand(J).isReg() &&
4709 HasSpecConstOperand(
Def->getOperand(J).getReg()))
4715 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
4716 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
4717 : SPIRV::OpConstantComposite;
4718 unsigned ContinuedOpc = HasSpecConst
4719 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
4720 : SPIRV::OpConstantCompositeContinuedINTEL;
4721 MachineIRBuilder MIR(
I);
4723 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
4725 for (
auto *Instr : Instructions) {
4726 Instr->setDebugLoc(
I.getDebugLoc());
4731 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4738 case Intrinsic::spv_assign_name: {
4739 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4740 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4741 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4742 i <
I.getNumExplicitOperands(); ++i) {
4743 MIB.
addImm(
I.getOperand(i).getImm());
4748 case Intrinsic::spv_switch: {
4749 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4750 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4751 if (
I.getOperand(i).isReg())
4752 MIB.
addReg(
I.getOperand(i).getReg());
4753 else if (
I.getOperand(i).isCImm())
4754 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4755 else if (
I.getOperand(i).isMBB())
4756 MIB.
addMBB(
I.getOperand(i).getMBB());
4763 case Intrinsic::spv_loop_merge: {
4764 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4765 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4766 if (
I.getOperand(i).isMBB())
4767 MIB.
addMBB(
I.getOperand(i).getMBB());
4774 case Intrinsic::spv_loop_control_intel: {
4776 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4777 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4782 case Intrinsic::spv_selection_merge: {
4784 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4785 assert(
I.getOperand(1).isMBB() &&
4786 "operand 1 to spv_selection_merge must be a basic block");
4787 MIB.
addMBB(
I.getOperand(1).getMBB());
4788 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4792 case Intrinsic::spv_cmpxchg:
4793 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4794 case Intrinsic::spv_unreachable:
4795 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4798 case Intrinsic::spv_abort:
4799 return selectAbort(
I);
4800 case Intrinsic::spv_alloca:
4801 return selectFrameIndex(ResVReg, ResType,
I);
4802 case Intrinsic::spv_alloca_array:
4803 return selectAllocaArray(ResVReg, ResType,
I);
4804 case Intrinsic::spv_assume:
4806 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4807 .
addUse(
I.getOperand(1).getReg())
4812 case Intrinsic::spv_expect:
4814 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4817 .
addUse(
I.getOperand(2).getReg())
4818 .
addUse(
I.getOperand(3).getReg())
4823 case Intrinsic::arithmetic_fence:
4824 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4825 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4828 .
addUse(
I.getOperand(2).getReg())
4832 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4834 case Intrinsic::spv_thread_id:
4840 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4842 case Intrinsic::spv_thread_id_in_group:
4848 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4850 case Intrinsic::spv_group_id:
4856 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4858 case Intrinsic::spv_flattened_thread_id_in_group:
4865 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4867 case Intrinsic::spv_workgroup_size:
4868 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4870 case Intrinsic::spv_global_size:
4871 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4873 case Intrinsic::spv_global_offset:
4874 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4876 case Intrinsic::spv_num_workgroups:
4877 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4879 case Intrinsic::spv_subgroup_size:
4880 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4882 case Intrinsic::spv_num_subgroups:
4883 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4885 case Intrinsic::spv_subgroup_id:
4886 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4887 case Intrinsic::spv_subgroup_local_invocation_id:
4888 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4889 ResVReg, ResType,
I);
4890 case Intrinsic::spv_subgroup_max_size:
4891 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4893 case Intrinsic::spv_fdot:
4894 return selectFloatDot(ResVReg, ResType,
I);
4895 case Intrinsic::spv_udot:
4896 case Intrinsic::spv_sdot:
4897 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4899 return selectIntegerDot(ResVReg, ResType,
I,
4900 IID == Intrinsic::spv_sdot);
4901 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4902 case Intrinsic::spv_dot4add_i8packed:
4903 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4905 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4906 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4907 case Intrinsic::spv_dot4add_u8packed:
4908 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4910 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4911 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4912 case Intrinsic::spv_all:
4913 return selectAll(ResVReg, ResType,
I);
4914 case Intrinsic::spv_any:
4915 return selectAny(ResVReg, ResType,
I);
4916 case Intrinsic::spv_cross:
4917 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4918 case Intrinsic::spv_distance:
4919 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4920 case Intrinsic::spv_lerp:
4921 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4922 case Intrinsic::spv_length:
4923 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4924 case Intrinsic::spv_degrees:
4925 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4926 case Intrinsic::spv_faceforward:
4927 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4928 case Intrinsic::spv_frac:
4929 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4930 case Intrinsic::spv_isinf:
4931 return selectOpIsInf(ResVReg, ResType,
I);
4932 case Intrinsic::spv_isnan:
4933 return selectOpIsNan(ResVReg, ResType,
I);
4934 case Intrinsic::spv_normalize:
4935 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4936 case Intrinsic::spv_refract:
4937 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4938 case Intrinsic::spv_reflect:
4939 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4940 case Intrinsic::spv_rsqrt:
4941 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4942 case Intrinsic::spv_sign:
4943 return selectSign(ResVReg, ResType,
I);
4944 case Intrinsic::spv_smoothstep:
4945 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4946 case Intrinsic::spv_firstbituhigh:
4947 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4948 case Intrinsic::spv_firstbitshigh:
4949 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4950 case Intrinsic::spv_firstbitlow:
4951 return selectFirstBitLow(ResVReg, ResType,
I);
4952 case Intrinsic::spv_all_memory_barrier:
4953 return selectBarrierInst(
I, SPIRV::Scope::Device,
4954 SPIRV::MemorySemantics::UniformMemory |
4955 SPIRV::MemorySemantics::ImageMemory |
4956 SPIRV::MemorySemantics::WorkgroupMemory,
4958 case Intrinsic::spv_all_memory_barrier_with_group_sync:
4959 return selectBarrierInst(
I, SPIRV::Scope::Device,
4960 SPIRV::MemorySemantics::UniformMemory |
4961 SPIRV::MemorySemantics::ImageMemory |
4962 SPIRV::MemorySemantics::WorkgroupMemory,
4964 case Intrinsic::spv_device_memory_barrier:
4965 return selectBarrierInst(
I, SPIRV::Scope::Device,
4966 SPIRV::MemorySemantics::UniformMemory |
4967 SPIRV::MemorySemantics::ImageMemory,
4969 case Intrinsic::spv_device_memory_barrier_with_group_sync:
4970 return selectBarrierInst(
I, SPIRV::Scope::Device,
4971 SPIRV::MemorySemantics::UniformMemory |
4972 SPIRV::MemorySemantics::ImageMemory,
4974 case Intrinsic::spv_group_memory_barrier:
4975 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4976 SPIRV::MemorySemantics::WorkgroupMemory,
4978 case Intrinsic::spv_group_memory_barrier_with_group_sync:
4979 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4980 SPIRV::MemorySemantics::WorkgroupMemory,
4982 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4983 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4984 SPIRV::StorageClass::StorageClass ResSC =
4988 "Generic storage class");
4989 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4997 case Intrinsic::spv_lifetime_start:
4998 case Intrinsic::spv_lifetime_end: {
4999 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5000 : SPIRV::OpLifetimeStop;
5001 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5002 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5011 case Intrinsic::spv_saturate:
5012 return selectSaturate(ResVReg, ResType,
I);
5013 case Intrinsic::spv_nclamp:
5014 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5015 case Intrinsic::spv_uclamp:
5016 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5017 case Intrinsic::spv_sclamp:
5018 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5019 case Intrinsic::spv_subgroup_prefix_bit_count:
5020 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5021 case Intrinsic::spv_wave_active_countbits:
5022 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5023 case Intrinsic::spv_wave_all_equal:
5024 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5025 case Intrinsic::spv_wave_all:
5026 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5027 case Intrinsic::spv_wave_any:
5028 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5029 case Intrinsic::spv_subgroup_ballot:
5030 return selectWaveOpInst(ResVReg, ResType,
I,
5031 SPIRV::OpGroupNonUniformBallot);
5032 case Intrinsic::spv_wave_is_first_lane:
5033 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5034 case Intrinsic::spv_wave_reduce_or:
5035 return selectWaveReduceOp(ResVReg, ResType,
I,
5036 SPIRV::OpGroupNonUniformBitwiseOr);
5037 case Intrinsic::spv_wave_reduce_xor:
5038 return selectWaveReduceOp(ResVReg, ResType,
I,
5039 SPIRV::OpGroupNonUniformBitwiseXor);
5040 case Intrinsic::spv_wave_reduce_and:
5041 return selectWaveReduceOp(ResVReg, ResType,
I,
5042 SPIRV::OpGroupNonUniformBitwiseAnd);
5043 case Intrinsic::spv_wave_reduce_umax:
5044 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5045 case Intrinsic::spv_wave_reduce_max:
5046 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5047 case Intrinsic::spv_wave_reduce_umin:
5048 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5049 case Intrinsic::spv_wave_reduce_min:
5050 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5051 case Intrinsic::spv_wave_reduce_sum:
5052 return selectWaveReduceSum(ResVReg, ResType,
I);
5053 case Intrinsic::spv_wave_product:
5054 return selectWaveReduceProduct(ResVReg, ResType,
I);
5055 case Intrinsic::spv_wave_readlane:
5056 return selectWaveOpInst(ResVReg, ResType,
I,
5057 SPIRV::OpGroupNonUniformShuffle);
5058 case Intrinsic::spv_wave_prefix_sum:
5059 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5060 case Intrinsic::spv_wave_prefix_product:
5061 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5062 case Intrinsic::spv_quad_read_across_x: {
5063 return selectQuadSwap(ResVReg, ResType,
I, 0);
5065 case Intrinsic::spv_quad_read_across_y: {
5066 return selectQuadSwap(ResVReg, ResType,
I, 1);
5068 case Intrinsic::spv_step:
5069 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5070 case Intrinsic::spv_radians:
5071 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5075 case Intrinsic::instrprof_increment:
5076 case Intrinsic::instrprof_increment_step:
5077 case Intrinsic::instrprof_value_profile:
5080 case Intrinsic::spv_value_md:
5082 case Intrinsic::spv_resource_handlefrombinding: {
5083 return selectHandleFromBinding(ResVReg, ResType,
I);
5085 case Intrinsic::spv_resource_counterhandlefrombinding:
5086 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5087 case Intrinsic::spv_resource_updatecounter:
5088 return selectUpdateCounter(ResVReg, ResType,
I);
5089 case Intrinsic::spv_resource_store_typedbuffer: {
5090 return selectImageWriteIntrinsic(
I);
5092 case Intrinsic::spv_resource_load_typedbuffer: {
5093 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5095 case Intrinsic::spv_resource_load_level: {
5096 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5098 case Intrinsic::spv_resource_getdimensions_x:
5099 case Intrinsic::spv_resource_getdimensions_xy:
5100 case Intrinsic::spv_resource_getdimensions_xyz: {
5101 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5103 case Intrinsic::spv_resource_getdimensions_levels_x:
5104 case Intrinsic::spv_resource_getdimensions_levels_xy:
5105 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5106 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5108 case Intrinsic::spv_resource_getdimensions_ms_xy:
5109 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5110 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5112 case Intrinsic::spv_resource_calculate_lod:
5113 case Intrinsic::spv_resource_calculate_lod_unclamped:
5114 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5115 case Intrinsic::spv_resource_sample:
5116 case Intrinsic::spv_resource_sample_clamp:
5117 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5118 case Intrinsic::spv_resource_samplebias:
5119 case Intrinsic::spv_resource_samplebias_clamp:
5120 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5121 case Intrinsic::spv_resource_samplegrad:
5122 case Intrinsic::spv_resource_samplegrad_clamp:
5123 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5124 case Intrinsic::spv_resource_samplelevel:
5125 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5126 case Intrinsic::spv_resource_samplecmp:
5127 case Intrinsic::spv_resource_samplecmp_clamp:
5128 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5129 case Intrinsic::spv_resource_samplecmplevelzero:
5130 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5131 case Intrinsic::spv_resource_gather:
5132 case Intrinsic::spv_resource_gather_cmp:
5133 return selectGatherIntrinsic(ResVReg, ResType,
I);
5134 case Intrinsic::spv_resource_getbasepointer:
5135 case Intrinsic::spv_resource_getpointer: {
5136 return selectResourceGetPointer(ResVReg, ResType,
I);
5138 case Intrinsic::spv_pushconstant_getpointer: {
5139 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5141 case Intrinsic::spv_discard: {
5142 return selectDiscard(ResVReg, ResType,
I);
5144 case Intrinsic::spv_resource_nonuniformindex: {
5145 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5147 case Intrinsic::spv_unpackhalf2x16: {
5148 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5150 case Intrinsic::spv_packhalf2x16: {
5151 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5153 case Intrinsic::spv_ddx:
5154 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5155 case Intrinsic::spv_ddy:
5156 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5157 case Intrinsic::spv_ddx_coarse:
5158 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5159 case Intrinsic::spv_ddy_coarse:
5160 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5161 case Intrinsic::spv_ddx_fine:
5162 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5163 case Intrinsic::spv_ddy_fine:
5164 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5165 case Intrinsic::spv_fwidth:
5166 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5167 case Intrinsic::spv_masked_gather:
5168 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5169 return selectMaskedGather(ResVReg, ResType,
I);
5170 return diagnoseUnsupported(
5171 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5172 case Intrinsic::spv_masked_scatter:
5173 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5174 return selectMaskedScatter(
I);
5175 return diagnoseUnsupported(
5176 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5177 case Intrinsic::returnaddress:
5178 case Intrinsic::frameaddress: {
5180 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5187 std::string DiagMsg;
5188 raw_string_ostream OS(DiagMsg);
5190 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
5197bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5198 SPIRVTypeInst ResType,
5199 MachineInstr &
I)
const {
5202 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5209bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5210 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5212 assert(Intr.getIntrinsicID() ==
5213 Intrinsic::spv_resource_counterhandlefrombinding);
5216 Register MainHandleReg = Intr.getOperand(2).getReg();
5218 assert(MainHandleDef->getIntrinsicID() ==
5219 Intrinsic::spv_resource_handlefrombinding);
5223 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5224 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5225 std::string CounterName =
5230 MachineIRBuilder MIRBuilder(
I);
5232 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5234 ArraySize, IndexReg, CounterName, MIRBuilder);
5236 return BuildCOPY(ResVReg, CounterVarReg,
I);
5239bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5240 SPIRVTypeInst ResType,
5241 MachineInstr &
I)
const {
5243 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5245 Register CounterHandleReg = Intr.getOperand(2).getReg();
5246 Register IncrReg = Intr.getOperand(3).getReg();
5253 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5254 assert(CounterVarPointeeType &&
5255 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5256 "Counter variable must be a struct");
5258 SPIRV::StorageClass::StorageBuffer &&
5259 "Counter variable must be in the storage buffer storage class");
5261 "Counter variable must have exactly 1 member in the struct");
5262 const SPIRVTypeInst MemberType =
5265 "Counter variable struct must have a single i32 member");
5269 MachineIRBuilder MIRBuilder(
I);
5271 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5274 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5280 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5283 .
addUse(CounterHandleReg)
5290 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5293 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5296 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5305 return BuildCOPY(ResVReg, AtomicRes,
I);
5313 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5321bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5322 SPIRVTypeInst ResType,
5323 MachineInstr &
I)
const {
5331 Register ImageReg =
I.getOperand(2).getReg();
5339 Register IdxReg =
I.getOperand(3).getReg();
5341 MachineInstr &Pos =
I;
5343 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5347bool SPIRVInstructionSelector::generateSampleImage(
5350 DebugLoc Loc, MachineInstr &Pos)
const {
5361 if (!loadHandleBeforePosition(NewSamplerReg,
5367 MachineIRBuilder MIRBuilder(Pos);
5380 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5381 ImOps.Lod.has_value();
5382 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5383 : SPIRV::OpImageSampleImplicitLod;
5385 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5386 : SPIRV::OpImageSampleDrefImplicitLod;
5395 MIB.
addUse(*ImOps.Compare);
5397 uint32_t ImageOperands = 0;
5399 ImageOperands |= SPIRV::ImageOperand::Bias;
5401 ImageOperands |= SPIRV::ImageOperand::Lod;
5402 if (ImOps.GradX && ImOps.GradY)
5403 ImageOperands |= SPIRV::ImageOperand::Grad;
5404 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5406 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5409 "Non-constant offsets are not supported in sample instructions.");
5413 ImageOperands |= SPIRV::ImageOperand::MinLod;
5415 if (ImageOperands != 0) {
5416 MIB.
addImm(ImageOperands);
5417 if (ImageOperands & SPIRV::ImageOperand::Bias)
5419 if (ImageOperands & SPIRV::ImageOperand::Lod)
5421 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5422 MIB.
addUse(*ImOps.GradX);
5423 MIB.
addUse(*ImOps.GradY);
5426 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5427 MIB.
addUse(*ImOps.Offset);
5428 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5429 MIB.
addUse(*ImOps.MinLod);
5436bool SPIRVInstructionSelector::selectImageQuerySize(
5438 std::optional<Register> LodReg)
const {
5440 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5443 "ImageReg is not an image type.");
5445 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5447 unsigned NumComponents = 0;
5449 case SPIRV::Dim::DIM_1D:
5450 case SPIRV::Dim::DIM_Buffer:
5451 NumComponents =
IsArray ? 2 : 1;
5453 case SPIRV::Dim::DIM_2D:
5454 case SPIRV::Dim::DIM_Cube:
5455 case SPIRV::Dim::DIM_Rect:
5456 NumComponents =
IsArray ? 3 : 2;
5458 case SPIRV::Dim::DIM_3D:
5462 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5467 SPIRVTypeInst ResType =
5472 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5482bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5483 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5484 Register ImageReg =
I.getOperand(2).getReg();
5491 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5494bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5495 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5496 Register ImageReg =
I.getOperand(2).getReg();
5505 Register LodReg =
I.getOperand(3).getReg();
5508 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5510 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5517 TII.get(SPIRV::OpImageQueryLevels))
5524 TII.get(SPIRV::OpCompositeConstruct))
5534bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5535 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5536 Register ImageReg =
I.getOperand(2).getReg();
5547 "OpImageQuerySamples requires a multisampled image");
5549 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5557 TII.get(SPIRV::OpImageQuerySamples))
5564 TII.get(SPIRV::OpCompositeConstruct))
5574bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5575 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5576 Register ImageReg =
I.getOperand(2).getReg();
5577 Register SamplerReg =
I.getOperand(3).getReg();
5578 Register CoordinateReg =
I.getOperand(4).getReg();
5594 if (!loadHandleBeforePosition(
5599 MachineIRBuilder MIRBuilder(
I);
5605 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5615 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
5622 unsigned ExtractedIndex =
5624 Intrinsic::spv_resource_calculate_lod_unclamped
5628 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5629 TII.get(SPIRV::OpCompositeExtract))
5639bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
5640 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5641 Register ImageReg =
I.getOperand(2).getReg();
5642 Register SamplerReg =
I.getOperand(3).getReg();
5643 Register CoordinateReg =
I.getOperand(4).getReg();
5644 ImageOperands ImOps;
5645 if (
I.getNumOperands() > 5)
5646 ImOps.Offset =
I.getOperand(5).getReg();
5647 if (
I.getNumOperands() > 6)
5648 ImOps.MinLod =
I.getOperand(6).getReg();
5649 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5650 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5653bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5654 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5655 Register ImageReg =
I.getOperand(2).getReg();
5656 Register SamplerReg =
I.getOperand(3).getReg();
5657 Register CoordinateReg =
I.getOperand(4).getReg();
5658 ImageOperands ImOps;
5659 ImOps.Bias =
I.getOperand(5).getReg();
5660 if (
I.getNumOperands() > 6)
5661 ImOps.Offset =
I.getOperand(6).getReg();
5662 if (
I.getNumOperands() > 7)
5663 ImOps.MinLod =
I.getOperand(7).getReg();
5664 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5665 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5668bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
5669 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5670 Register ImageReg =
I.getOperand(2).getReg();
5671 Register SamplerReg =
I.getOperand(3).getReg();
5672 Register CoordinateReg =
I.getOperand(4).getReg();
5673 ImageOperands ImOps;
5674 ImOps.GradX =
I.getOperand(5).getReg();
5675 ImOps.GradY =
I.getOperand(6).getReg();
5676 if (
I.getNumOperands() > 7)
5677 ImOps.Offset =
I.getOperand(7).getReg();
5678 if (
I.getNumOperands() > 8)
5679 ImOps.MinLod =
I.getOperand(8).getReg();
5680 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5681 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5684bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
5685 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5686 Register ImageReg =
I.getOperand(2).getReg();
5687 Register SamplerReg =
I.getOperand(3).getReg();
5688 Register CoordinateReg =
I.getOperand(4).getReg();
5689 ImageOperands ImOps;
5690 ImOps.Lod =
I.getOperand(5).getReg();
5691 if (
I.getNumOperands() > 6)
5692 ImOps.Offset =
I.getOperand(6).getReg();
5693 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5694 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5697bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
5698 SPIRVTypeInst ResType,
5699 MachineInstr &
I)
const {
5700 Register ImageReg =
I.getOperand(2).getReg();
5701 Register SamplerReg =
I.getOperand(3).getReg();
5702 Register CoordinateReg =
I.getOperand(4).getReg();
5703 ImageOperands ImOps;
5704 ImOps.Compare =
I.getOperand(5).getReg();
5705 if (
I.getNumOperands() > 6)
5706 ImOps.Offset =
I.getOperand(6).getReg();
5707 if (
I.getNumOperands() > 7)
5708 ImOps.MinLod =
I.getOperand(7).getReg();
5709 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5710 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5713bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
5714 SPIRVTypeInst ResType,
5715 MachineInstr &
I)
const {
5716 Register ImageReg =
I.getOperand(2).getReg();
5717 Register CoordinateReg =
I.getOperand(3).getReg();
5718 Register LodReg =
I.getOperand(4).getReg();
5720 ImageOperands ImOps;
5722 if (
I.getNumOperands() > 5)
5723 ImOps.Offset =
I.getOperand(5).getReg();
5735 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
5736 I.getDebugLoc(),
I, &ImOps);
5739bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
5740 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5741 Register ImageReg =
I.getOperand(2).getReg();
5742 Register SamplerReg =
I.getOperand(3).getReg();
5743 Register CoordinateReg =
I.getOperand(4).getReg();
5744 ImageOperands ImOps;
5745 ImOps.Compare =
I.getOperand(5).getReg();
5746 if (
I.getNumOperands() > 6)
5747 ImOps.Offset =
I.getOperand(6).getReg();
5750 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5751 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5754bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
5755 SPIRVTypeInst ResType,
5756 MachineInstr &
I)
const {
5757 Register ImageReg =
I.getOperand(2).getReg();
5758 Register SamplerReg =
I.getOperand(3).getReg();
5759 Register CoordinateReg =
I.getOperand(4).getReg();
5762 "ImageReg is not an image type.");
5767 ComponentOrCompareReg =
I.getOperand(5).getReg();
5768 OffsetReg =
I.getOperand(6).getReg();
5771 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
5775 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5776 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
5777 Dim != SPIRV::Dim::DIM_Rect) {
5779 "Gather operations are only supported for 2D, Cube, and Rect images.");
5786 if (!loadHandleBeforePosition(
5791 MachineIRBuilder MIRBuilder(
I);
5792 SPIRVTypeInst SampledImageType =
5797 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5805 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
5807 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
5809 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5814 .
addUse(ComponentOrCompareReg);
5816 uint32_t ImageOperands = 0;
5817 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
5818 if (Dim == SPIRV::Dim::DIM_Cube) {
5820 "Gather operations with offset are not supported for Cube images.");
5824 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5826 ImageOperands |= SPIRV::ImageOperand::Offset;
5830 if (ImageOperands != 0) {
5831 MIB.
addImm(ImageOperands);
5833 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5841bool SPIRVInstructionSelector::generateImageReadOrFetch(
5844 const ImageOperands *ImOps)
const {
5847 "ImageReg is not an image type.");
5849 bool IsSignedInteger =
5854 bool IsFetch = (SampledOp.getImm() == 1);
5856 auto AddOperands = [&](MachineInstrBuilder &MIB) {
5857 uint32_t ImageOperandsMask = 0;
5858 if (IsSignedInteger)
5859 ImageOperandsMask |= 0x1000;
5861 if (IsFetch && ImOps) {
5863 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
5864 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
5866 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
5868 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5872 if (ImageOperandsMask != 0) {
5873 MIB.
addImm(ImageOperandsMask);
5874 if (IsFetch && ImOps) {
5877 if (ImOps->Offset &&
5878 (ImageOperandsMask &
5879 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5880 MIB.
addUse(*ImOps->Offset);
5886 if (ResultSize == 4) {
5889 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5896 BMI.constrainAllUses(
TII,
TRI, RBI);
5900 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5904 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5910 BMI.constrainAllUses(
TII,
TRI, RBI);
5912 if (ResultSize == 1) {
5921 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
5924bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
5925 SPIRVTypeInst ResType,
5926 MachineInstr &
I)
const {
5927 Register ResourcePtr =
I.getOperand(2).getReg();
5929 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
5938 MachineIRBuilder MIRBuilder(
I);
5943 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5949 if (
I.getNumExplicitOperands() > 3) {
5950 Register IndexReg =
I.getOperand(3).getReg();
5957bool SPIRVInstructionSelector::selectPushConstantGetPointer(
5958 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5963bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
5964 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5965 Register ObjReg =
I.getOperand(2).getReg();
5966 if (!BuildCOPY(ResVReg, ObjReg,
I))
5976 decorateUsesAsNonUniform(ResVReg);
5980void SPIRVInstructionSelector::decorateUsesAsNonUniform(
5983 while (WorkList.
size() > 0) {
5987 bool IsDecorated =
false;
5989 if (
Use.getOpcode() == SPIRV::OpDecorate &&
5990 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
5996 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
5998 if (ResultReg == CurrentReg)
6006 SPIRV::Decoration::NonUniformEXT, {});
6011bool SPIRVInstructionSelector::extractSubvector(
6013 MachineInstr &InsertionPoint)
const {
6015 [[maybe_unused]] uint64_t InputSize =
6018 assert(InputSize > 1 &&
"The input must be a vector.");
6019 assert(ResultSize > 1 &&
"The result must be a vector.");
6020 assert(ResultSize < InputSize &&
6021 "Cannot extract more element than there are in the input.");
6024 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
6025 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6028 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6037 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6039 TII.get(SPIRV::OpCompositeConstruct))
6043 for (
Register ComponentReg : ComponentRegisters)
6044 MIB.
addUse(ComponentReg);
6049bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6050 MachineInstr &
I)
const {
6057 Register ImageReg =
I.getOperand(1).getReg();
6065 Register CoordinateReg =
I.getOperand(2).getReg();
6066 Register DataReg =
I.getOperand(3).getReg();
6069 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6077Register SPIRVInstructionSelector::buildPointerToResource(
6078 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6079 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6080 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6082 if (ArraySize == 1) {
6083 SPIRVTypeInst PtrType =
6086 "SpirvResType did not have an explicit layout.");
6091 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6092 SPIRVTypeInst VarPointerType =
6095 VarPointerType, Set,
Binding, Name, MIRBuilder);
6097 SPIRVTypeInst ResPointerType =
6110bool SPIRVInstructionSelector::selectFirstBitSet16(
6111 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6112 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6114 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6118 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6121bool SPIRVInstructionSelector::selectFirstBitSet32(
6123 unsigned BitSetOpcode)
const {
6124 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6127 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6134bool SPIRVInstructionSelector::selectFirstBitSet64(
6136 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6149 if (ComponentCount > 2) {
6150 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6152 unsigned Opcode) ->
bool {
6153 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6157 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6161 MachineIRBuilder MIRBuilder(
I);
6163 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6167 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6173 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6180 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6183 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6184 SPIRV::OpVectorExtractDynamic))
6186 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6187 SPIRV::OpVectorExtractDynamic))
6191 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6192 TII.get(SPIRV::OpVectorShuffle))
6200 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6206 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6207 TII.get(SPIRV::OpVectorShuffle))
6215 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6235 SelectOp = SPIRV::OpSelectSISCond;
6236 AddOp = SPIRV::OpIAddS;
6244 SelectOp = SPIRV::OpSelectVIVCond;
6245 AddOp = SPIRV::OpIAddV;
6251 Register RegSecondaryOffset = Reg0;
6255 if (SwapPrimarySide) {
6256 PrimaryReg = LowReg;
6257 SecondaryReg = HighReg;
6258 RegPrimaryOffset = Reg0;
6259 RegSecondaryOffset = Reg32;
6264 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6265 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6270 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6271 SPIRV::OpINotEqual))
6278 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6279 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6284 if (SwapPrimarySide) {
6286 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6287 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6298 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6299 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6304 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6305 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6308 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6312bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6313 SPIRVTypeInst ResType,
6315 bool IsSigned)
const {
6317 Register OpReg =
I.getOperand(2).getReg();
6320 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6321 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6325 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6327 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6329 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6333 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6337bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6338 SPIRVTypeInst ResType,
6339 MachineInstr &
I)
const {
6341 Register OpReg =
I.getOperand(2).getReg();
6346 unsigned ExtendOpcode = SPIRV::OpUConvert;
6347 unsigned BitSetOpcode = GL::FindILsb;
6351 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6353 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6355 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6362bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6363 SPIRVTypeInst ResType,
6364 MachineInstr &
I)
const {
6368 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6371 .
addUse(
I.getOperand(2).getReg())
6374 unsigned Alignment =
I.getOperand(3).getImm();
6388 while (!Worklist.
empty()) {
6390 switch (
T->getOpcode()) {
6391 case SPIRV::OpTypeInt:
6392 case SPIRV::OpTypeFloat:
6393 case SPIRV::OpTypePointer:
6395 case SPIRV::OpTypeVector:
6396 case SPIRV::OpTypeMatrix:
6397 case SPIRV::OpTypeArray: {
6398 Register OperandReg =
T->getOperand(1).getReg();
6402 case SPIRV::OpTypeStruct:
6403 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6404 Register OperandReg =
T->getOperand(Idx).getReg();
6416bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6417 assert(
I.getNumExplicitOperands() == 2);
6419 Register MsgReg =
I.getOperand(1).getReg();
6421 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6424 return diagnoseUnsupported(
6426 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6427 "scalar, pointer, vector, matrix, or aggregate of such types)");
6430 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6437bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6446 uint32_t MsgVal = ~0
u;
6447 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6448 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6451 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6454 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6461bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6462 SPIRVTypeInst ResType,
6463 MachineInstr &
I)
const {
6467 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6470 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6473 unsigned Alignment =
I.getOperand(2).getImm();
6480bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6485 const MachineInstr *PrevI =
I.getPrevNode();
6487 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6491 .
addMBB(
I.getOperand(0).getMBB())
6496 .
addMBB(
I.getOperand(0).getMBB())
6501bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6512 const MachineInstr *NextI =
I.getNextNode();
6514 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6520 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6522 .
addUse(
I.getOperand(0).getReg())
6523 .
addMBB(
I.getOperand(1).getMBB())
6529bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6530 MachineInstr &
I)
const {
6532 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6534 const unsigned NumOps =
I.getNumOperands();
6535 for (
unsigned i = 1; i <
NumOps; i += 2) {
6536 MIB.
addUse(
I.getOperand(i + 0).getReg());
6537 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6543bool SPIRVInstructionSelector::selectGlobalValue(
6544 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6546 MachineIRBuilder MIRBuilder(
I);
6547 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6550 std::string GlobalIdent;
6552 unsigned &
ID = UnnamedGlobalIDs[GV];
6554 ID = UnnamedGlobalIDs.
size();
6555 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6581 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6588 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6593 MachineInstrBuilder MIB1 =
6594 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6597 MachineInstrBuilder MIB2 =
6599 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
6603 GR.
add(ConstVal, MIB2);
6611 MachineInstrBuilder MIB3 =
6612 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
6615 GR.
add(ConstVal, MIB3);
6619 assert(NewReg != ResVReg);
6620 return BuildCOPY(ResVReg, NewReg,
I);
6630 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
6633 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
6639 SPIRVTypeInst ResType =
6643 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6648 if (
GlobalVar->isExternallyInitialized() &&
6649 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6650 constexpr unsigned ReadWriteINTEL = 3u;
6653 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6659bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6660 SPIRVTypeInst ResType,
6661 MachineInstr &
I)
const {
6663 return selectExtInst(ResVReg, ResType,
I, CL::log10);
6671 MachineIRBuilder MIRBuilder(
I);
6676 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6679 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6681 .
add(
I.getOperand(1))
6686 ResType->
getOpcode() == SPIRV::OpTypeFloat);
6696 APFloat::rmNearestTiesToEven, &LosesInfo);
6700 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
6701 ? SPIRV::OpVectorTimesScalar
6712bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
6713 SPIRVTypeInst ResType,
6714 MachineInstr &
I)
const {
6717 return selectExtInst(ResVReg, ResType,
I, CL::pown);
6723 Register ExpReg =
I.getOperand(2).getReg();
6725 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
6726 SPIRV::OpConvertSToF))
6728 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
6735bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
6736 SPIRVTypeInst ResType,
6737 MachineInstr &
I)
const {
6753 MachineIRBuilder MIRBuilder(
I);
6756 ResType, MIRBuilder, SPIRV::StorageClass::Function);
6769 MachineBasicBlock &EntryBB =
I.getMF()->
front();
6771 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
6774 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
6780 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6783 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
6786 .
add(
I.getOperand(
I.getNumExplicitDefs()))
6790 Register IntegralPartReg =
I.getOperand(1).getReg();
6791 if (IntegralPartReg.
isValid()) {
6793 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6804 assert(
false &&
"GLSL::Modf is deprecated.");
6815bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
6816 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6817 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6818 MachineIRBuilder MIRBuilder(
I);
6819 const SPIRVTypeInst Vec3Ty =
6822 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
6834 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6838 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6844 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6851 assert(
I.getOperand(2).isReg());
6852 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
6856 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6867bool SPIRVInstructionSelector::loadBuiltinInputID(
6868 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6869 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6870 MachineIRBuilder MIRBuilder(
I);
6872 ResType, MIRBuilder, SPIRV::StorageClass::Input);
6887 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6891 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6900SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
6901 MachineInstr &
I)
const {
6902 MachineIRBuilder MIRBuilder(
I);
6903 if (
Type->getOpcode() != SPIRV::OpTypeVector)
6913bool SPIRVInstructionSelector::loadHandleBeforePosition(
6914 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
6915 MachineInstr &Pos)
const {
6918 Intrinsic::spv_resource_handlefrombinding);
6926 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
6927 MachineIRBuilder MIRBuilder(HandleDef);
6928 SPIRVTypeInst VarType = ResType;
6929 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
6931 if (IsStructuredBuffer) {
6936 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
6938 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
6941 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
6942 ArraySize, IndexReg, Name, MIRBuilder);
6946 uint32_t LoadOpcode =
6947 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
6957void SPIRVInstructionSelector::errorIfInstrOutsideShader(
6958 MachineInstr &
I)
const {
6960 std::string DiagMsg;
6961 raw_string_ostream OS(DiagMsg);
6962 I.print(OS,
true,
false,
false,
false);
6963 DiagMsg +=
" is only supported in shaders.\n";
6969InstructionSelector *
6973 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...