34#include "llvm/IR/IntrinsicsSPIRV.h"
38#define DEBUG_TYPE "spirv-isel"
45 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
50 std::optional<Register> Bias;
51 std::optional<Register>
Offset;
52 std::optional<Register> MinLod;
53 std::optional<Register> GradX;
54 std::optional<Register> GradY;
55 std::optional<Register> Lod;
56 std::optional<Register> Compare;
59llvm::SPIRV::SelectionControl::SelectionControl
60getSelectionOperandForImm(
int Imm) {
62 return SPIRV::SelectionControl::Flatten;
64 return SPIRV::SelectionControl::DontFlatten;
66 return SPIRV::SelectionControl::None;
70#define GET_GLOBALISEL_PREDICATE_BITSET
71#include "SPIRVGenGlobalISel.inc"
72#undef GET_GLOBALISEL_PREDICATE_BITSET
99#define GET_GLOBALISEL_PREDICATES_DECL
100#include "SPIRVGenGlobalISel.inc"
101#undef GET_GLOBALISEL_PREDICATES_DECL
103#define GET_GLOBALISEL_TEMPORARIES_DECL
104#include "SPIRVGenGlobalISel.inc"
105#undef GET_GLOBALISEL_TEMPORARIES_DECL
129 unsigned BitSetOpcode)
const;
133 unsigned BitSetOpcode)
const;
137 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
141 unsigned BitSetOpcode,
142 bool SwapPrimarySide)
const;
149 unsigned Opcode)
const;
152 unsigned Opcode)
const;
171 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
182 unsigned OpType)
const;
237 template <
bool Signed>
240 template <
bool Signed>
247 template <
typename PickOpcodeFn>
250 PickOpcodeFn &&PickOpcode)
const;
267 template <
typename PickOpcodeFn>
270 PickOpcodeFn &&PickOpcode)
const;
288 bool IsSigned)
const;
290 bool IsSigned,
unsigned Opcode)
const;
292 bool IsSigned)
const;
298 bool IsSigned)
const;
337 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
338 bool useMISrc =
true,
340 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
341 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
342 bool useMISrc =
true,
344 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
345 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
346 bool setMIFlags =
true,
bool useMISrc =
true,
348 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
349 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
350 bool useMISrc =
true,
353 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
354 MachineInstr &
I)
const;
356 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
357 MachineInstr &
I)
const;
359 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
360 MachineInstr &
I)
const;
362 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
363 MachineInstr &
I,
unsigned Opcode)
const;
365 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
366 bool WithGroupSync)
const;
368 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
369 MachineInstr &
I)
const;
371 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
372 MachineInstr &
I)
const;
376 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
377 MachineInstr &
I)
const;
379 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
380 MachineInstr &
I)
const;
382 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
383 MachineInstr &
I)
const;
384 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
385 MachineInstr &
I)
const;
386 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
387 MachineInstr &
I)
const;
388 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
389 MachineInstr &
I)
const;
390 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
391 MachineInstr &
I)
const;
392 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
393 MachineInstr &
I)
const;
394 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
395 MachineInstr &
I)
const;
396 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
397 MachineInstr &
I)
const;
398 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
399 SPIRVTypeInst ResType,
400 MachineInstr &
I)
const;
401 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
402 MachineInstr &
I)
const;
403 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
404 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
405 MachineInstr &
I)
const;
406 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
407 MachineInstr &
I)
const;
408 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
409 MachineInstr &
I)
const;
410 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
411 MachineInstr &
I)
const;
412 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
414 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
415 MachineInstr &
I)
const;
416 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
417 MachineInstr &
I)
const;
418 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
419 MachineInstr &
I)
const;
420 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
421 MachineInstr &
I,
const unsigned DPdOpCode)
const;
423 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
424 SPIRVTypeInst ResType =
nullptr)
const;
426 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
427 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
428 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
430 MachineInstr &
I)
const;
431 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
433 bool wrapIntoSpecConstantOp(MachineInstr &
I,
436 Register getUcharPtrTypeReg(MachineInstr &
I,
437 SPIRV::StorageClass::StorageClass SC)
const;
438 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
440 uint32_t Opcode)
const;
441 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
442 SPIRVTypeInst SrcPtrTy)
const;
443 Register buildPointerToResource(SPIRVTypeInst ResType,
444 SPIRV::StorageClass::StorageClass SC,
445 uint32_t Set, uint32_t
Binding,
446 uint32_t ArraySize,
Register IndexReg,
448 MachineIRBuilder MIRBuilder)
const;
449 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
450 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
451 Register &ReadReg, MachineInstr &InsertionPoint)
const;
452 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
455 const ImageOperands *ImOps =
nullptr)
const;
456 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
458 Register CoordinateReg,
const ImageOperands &ImOps,
461 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
462 Register ResVReg, SPIRVTypeInst ResType,
463 MachineInstr &
I)
const;
464 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
465 Register ResVReg, SPIRVTypeInst ResType,
466 MachineInstr &
I)
const;
467 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
468 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
469 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
470 void errorIfInstrOutsideShader(MachineInstr &
I)
const;
473bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
475 if (
TET->getTargetExtName() ==
"spirv.Image") {
478 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
479 return TET->getTypeParameter(0)->isIntegerTy();
483#define GET_GLOBALISEL_IMPL
484#include "SPIRVGenGlobalISel.inc"
485#undef GET_GLOBALISEL_IMPL
491 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
494#include
"SPIRVGenGlobalISel.inc"
497#include
"SPIRVGenGlobalISel.inc"
509 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
513void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
514 if (HasVRegsReset == &MF)
529 for (
const auto &
MBB : MF) {
530 for (
const auto &
MI :
MBB) {
533 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
537 LLT DstType = MRI.
getType(DstReg);
539 LLT SrcType = MRI.
getType(SrcReg);
540 if (DstType != SrcType)
545 if (DstRC != SrcRC && SrcRC)
557 while (!Stack.empty()) {
562 switch (
MI->getOpcode()) {
563 case TargetOpcode::G_INTRINSIC:
564 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
565 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
568 if (IntrID != Intrinsic::spv_const_composite &&
569 IntrID != Intrinsic::spv_undef)
573 case TargetOpcode::G_BUILD_VECTOR:
574 case TargetOpcode::G_SPLAT_VECTOR:
576 i < OpDef->getNumOperands(); i++) {
581 Stack.push_back(OpNestedDef);
584 case TargetOpcode::G_CONSTANT:
585 case TargetOpcode::G_FCONSTANT:
586 case TargetOpcode::G_IMPLICIT_DEF:
587 case SPIRV::OpConstantTrue:
588 case SPIRV::OpConstantFalse:
589 case SPIRV::OpConstantI:
590 case SPIRV::OpConstantF:
591 case SPIRV::OpConstantComposite:
592 case SPIRV::OpConstantCompositeContinuedINTEL:
593 case SPIRV::OpConstantSampler:
594 case SPIRV::OpConstantNull:
596 case SPIRV::OpConstantFunctionPointerINTEL:
623 case Intrinsic::spv_all:
624 case Intrinsic::spv_alloca:
625 case Intrinsic::spv_any:
626 case Intrinsic::spv_bitcast:
627 case Intrinsic::spv_const_composite:
628 case Intrinsic::spv_cross:
629 case Intrinsic::spv_degrees:
630 case Intrinsic::spv_distance:
631 case Intrinsic::spv_extractelt:
632 case Intrinsic::spv_extractv:
633 case Intrinsic::spv_faceforward:
634 case Intrinsic::spv_fdot:
635 case Intrinsic::spv_firstbitlow:
636 case Intrinsic::spv_firstbitshigh:
637 case Intrinsic::spv_firstbituhigh:
638 case Intrinsic::spv_frac:
639 case Intrinsic::spv_gep:
640 case Intrinsic::spv_global_offset:
641 case Intrinsic::spv_global_size:
642 case Intrinsic::spv_group_id:
643 case Intrinsic::spv_insertelt:
644 case Intrinsic::spv_insertv:
645 case Intrinsic::spv_isinf:
646 case Intrinsic::spv_isnan:
647 case Intrinsic::spv_lerp:
648 case Intrinsic::spv_length:
649 case Intrinsic::spv_normalize:
650 case Intrinsic::spv_num_subgroups:
651 case Intrinsic::spv_num_workgroups:
652 case Intrinsic::spv_ptrcast:
653 case Intrinsic::spv_radians:
654 case Intrinsic::spv_reflect:
655 case Intrinsic::spv_refract:
656 case Intrinsic::spv_resource_getpointer:
657 case Intrinsic::spv_resource_handlefrombinding:
658 case Intrinsic::spv_resource_handlefromimplicitbinding:
659 case Intrinsic::spv_resource_nonuniformindex:
660 case Intrinsic::spv_resource_sample:
661 case Intrinsic::spv_rsqrt:
662 case Intrinsic::spv_saturate:
663 case Intrinsic::spv_sdot:
664 case Intrinsic::spv_sign:
665 case Intrinsic::spv_smoothstep:
666 case Intrinsic::spv_step:
667 case Intrinsic::spv_subgroup_id:
668 case Intrinsic::spv_subgroup_local_invocation_id:
669 case Intrinsic::spv_subgroup_max_size:
670 case Intrinsic::spv_subgroup_size:
671 case Intrinsic::spv_thread_id:
672 case Intrinsic::spv_thread_id_in_group:
673 case Intrinsic::spv_udot:
674 case Intrinsic::spv_undef:
675 case Intrinsic::spv_value_md:
676 case Intrinsic::spv_workgroup_size:
688 case SPIRV::OpTypeVoid:
689 case SPIRV::OpTypeBool:
690 case SPIRV::OpTypeInt:
691 case SPIRV::OpTypeFloat:
692 case SPIRV::OpTypeVector:
693 case SPIRV::OpTypeMatrix:
694 case SPIRV::OpTypeImage:
695 case SPIRV::OpTypeSampler:
696 case SPIRV::OpTypeSampledImage:
697 case SPIRV::OpTypeArray:
698 case SPIRV::OpTypeRuntimeArray:
699 case SPIRV::OpTypeStruct:
700 case SPIRV::OpTypeOpaque:
701 case SPIRV::OpTypePointer:
702 case SPIRV::OpTypeFunction:
703 case SPIRV::OpTypeEvent:
704 case SPIRV::OpTypeDeviceEvent:
705 case SPIRV::OpTypeReserveId:
706 case SPIRV::OpTypeQueue:
707 case SPIRV::OpTypePipe:
708 case SPIRV::OpTypeForwardPointer:
709 case SPIRV::OpTypePipeStorage:
710 case SPIRV::OpTypeNamedBarrier:
711 case SPIRV::OpTypeAccelerationStructureNV:
712 case SPIRV::OpTypeCooperativeMatrixNV:
713 case SPIRV::OpTypeCooperativeMatrixKHR:
723 if (
MI.getNumDefs() == 0)
726 for (
const auto &MO :
MI.all_defs()) {
728 if (
Reg.isPhysical()) {
733 if (
UseMI.getOpcode() != SPIRV::OpName) {
740 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
741 MI.isLifetimeMarker()) {
744 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
755 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
756 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
759 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
764 if (
MI.mayStore() ||
MI.isCall() ||
765 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
766 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
767 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
778 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
785void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
787 for (
const auto &MO :
MI.all_defs()) {
791 SmallVector<MachineInstr *, 4> UselessOpNames;
794 "There is still a use of the dead function.");
797 for (MachineInstr *OpNameMI : UselessOpNames) {
799 OpNameMI->eraseFromParent();
804void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
807 removeOpNamesForDeadMI(
MI);
808 MI.eraseFromParent();
811bool SPIRVInstructionSelector::select(MachineInstr &
I) {
812 resetVRegsType(*
I.getParent()->getParent());
814 assert(
I.getParent() &&
"Instruction should be in a basic block!");
815 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
820 removeDeadInstruction(
I);
827 if (Opcode == SPIRV::ASSIGN_TYPE) {
828 Register DstReg =
I.getOperand(0).getReg();
829 Register SrcReg =
I.getOperand(1).getReg();
832 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
833 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
834 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
835 Register SelectDstReg =
Def->getOperand(0).getReg();
836 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
838 assert(SuccessToSelectSelect);
840 Def->eraseFromParent();
847 bool Res = selectImpl(
I, *CoverageInfo);
849 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
850 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
854 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
866 }
else if (
I.getNumDefs() == 1) {
878 removeDeadInstruction(
I);
883 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
884 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
890 bool HasDefs =
I.getNumDefs() > 0;
893 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
894 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
895 if (spvSelect(ResVReg, ResType,
I)) {
897 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
908 case TargetOpcode::G_CONSTANT:
909 case TargetOpcode::G_FCONSTANT:
916 MachineInstr &
I)
const {
919 if (DstRC != SrcRC && SrcRC)
921 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
928bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
929 SPIRVTypeInst ResType,
930 MachineInstr &
I)
const {
931 const unsigned Opcode =
I.getOpcode();
933 return selectImpl(
I, *CoverageInfo);
935 case TargetOpcode::G_CONSTANT:
936 case TargetOpcode::G_FCONSTANT:
937 return selectConst(ResVReg, ResType,
I);
938 case TargetOpcode::G_GLOBAL_VALUE:
939 return selectGlobalValue(ResVReg,
I);
940 case TargetOpcode::G_IMPLICIT_DEF:
941 return selectOpUndef(ResVReg, ResType,
I);
942 case TargetOpcode::G_FREEZE:
943 return selectFreeze(ResVReg, ResType,
I);
945 case TargetOpcode::G_INTRINSIC:
946 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
947 case TargetOpcode::G_INTRINSIC_CONVERGENT:
948 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
949 return selectIntrinsic(ResVReg, ResType,
I);
950 case TargetOpcode::G_BITREVERSE:
951 return selectBitreverse(ResVReg, ResType,
I);
953 case TargetOpcode::G_BUILD_VECTOR:
954 return selectBuildVector(ResVReg, ResType,
I);
955 case TargetOpcode::G_SPLAT_VECTOR:
956 return selectSplatVector(ResVReg, ResType,
I);
958 case TargetOpcode::G_SHUFFLE_VECTOR: {
959 MachineBasicBlock &BB = *
I.getParent();
960 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
963 .
addUse(
I.getOperand(1).getReg())
964 .
addUse(
I.getOperand(2).getReg());
965 for (
auto V :
I.getOperand(3).getShuffleMask())
970 case TargetOpcode::G_MEMMOVE:
971 case TargetOpcode::G_MEMCPY:
972 case TargetOpcode::G_MEMSET:
973 return selectMemOperation(ResVReg,
I);
975 case TargetOpcode::G_ICMP:
976 return selectICmp(ResVReg, ResType,
I);
977 case TargetOpcode::G_FCMP:
978 return selectFCmp(ResVReg, ResType,
I);
980 case TargetOpcode::G_FRAME_INDEX:
981 return selectFrameIndex(ResVReg, ResType,
I);
983 case TargetOpcode::G_LOAD:
984 return selectLoad(ResVReg, ResType,
I);
985 case TargetOpcode::G_STORE:
986 return selectStore(
I);
988 case TargetOpcode::G_BR:
989 return selectBranch(
I);
990 case TargetOpcode::G_BRCOND:
991 return selectBranchCond(
I);
993 case TargetOpcode::G_PHI:
994 return selectPhi(ResVReg,
I);
996 case TargetOpcode::G_FPTOSI:
997 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
998 case TargetOpcode::G_FPTOUI:
999 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1001 case TargetOpcode::G_FPTOSI_SAT:
1002 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1003 case TargetOpcode::G_FPTOUI_SAT:
1004 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1006 case TargetOpcode::G_SITOFP:
1007 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1008 case TargetOpcode::G_UITOFP:
1009 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1011 case TargetOpcode::G_CTPOP:
1012 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitCount);
1013 case TargetOpcode::G_SMIN:
1014 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1015 case TargetOpcode::G_UMIN:
1016 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1018 case TargetOpcode::G_SMAX:
1019 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1020 case TargetOpcode::G_UMAX:
1021 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1023 case TargetOpcode::G_SCMP:
1024 return selectSUCmp(ResVReg, ResType,
I,
true);
1025 case TargetOpcode::G_UCMP:
1026 return selectSUCmp(ResVReg, ResType,
I,
false);
1027 case TargetOpcode::G_LROUND:
1028 case TargetOpcode::G_LLROUND: {
1031 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1033 regForLround, *(
I.getParent()->getParent()));
1035 CL::round, GL::Round,
false);
1037 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1044 case TargetOpcode::G_STRICT_FMA:
1045 case TargetOpcode::G_FMA: {
1048 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1051 .
addUse(
I.getOperand(1).getReg())
1052 .
addUse(
I.getOperand(2).getReg())
1053 .
addUse(
I.getOperand(3).getReg())
1058 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1061 case TargetOpcode::G_STRICT_FLDEXP:
1062 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1064 case TargetOpcode::G_FPOW:
1065 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1066 case TargetOpcode::G_FPOWI:
1067 return selectFpowi(ResVReg, ResType,
I);
1069 case TargetOpcode::G_FEXP:
1070 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1071 case TargetOpcode::G_FEXP2:
1072 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1073 case TargetOpcode::G_FEXP10:
1074 return selectExp10(ResVReg, ResType,
I);
1076 case TargetOpcode::G_FMODF:
1077 return selectModf(ResVReg, ResType,
I);
1078 case TargetOpcode::G_FSINCOS:
1079 return selectSincos(ResVReg, ResType,
I);
1081 case TargetOpcode::G_FLOG:
1082 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1083 case TargetOpcode::G_FLOG2:
1084 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1085 case TargetOpcode::G_FLOG10:
1086 return selectLog10(ResVReg, ResType,
I);
1088 case TargetOpcode::G_FABS:
1089 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1090 case TargetOpcode::G_ABS:
1091 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1093 case TargetOpcode::G_FMINNUM:
1094 case TargetOpcode::G_FMINIMUM:
1095 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1096 case TargetOpcode::G_FMAXNUM:
1097 case TargetOpcode::G_FMAXIMUM:
1098 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1100 case TargetOpcode::G_FCOPYSIGN:
1101 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1103 case TargetOpcode::G_FCEIL:
1104 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1105 case TargetOpcode::G_FFLOOR:
1106 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1108 case TargetOpcode::G_FCOS:
1109 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1110 case TargetOpcode::G_FSIN:
1111 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1112 case TargetOpcode::G_FTAN:
1113 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1114 case TargetOpcode::G_FACOS:
1115 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1116 case TargetOpcode::G_FASIN:
1117 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1118 case TargetOpcode::G_FATAN:
1119 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1120 case TargetOpcode::G_FATAN2:
1121 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1122 case TargetOpcode::G_FCOSH:
1123 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1124 case TargetOpcode::G_FSINH:
1125 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1126 case TargetOpcode::G_FTANH:
1127 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1129 case TargetOpcode::G_STRICT_FSQRT:
1130 case TargetOpcode::G_FSQRT:
1131 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1133 case TargetOpcode::G_CTTZ:
1134 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1135 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1136 case TargetOpcode::G_CTLZ:
1137 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1138 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1140 case TargetOpcode::G_INTRINSIC_ROUND:
1141 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1142 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1143 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1144 case TargetOpcode::G_INTRINSIC_TRUNC:
1145 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1146 case TargetOpcode::G_FRINT:
1147 case TargetOpcode::G_FNEARBYINT:
1148 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1150 case TargetOpcode::G_SMULH:
1151 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1152 case TargetOpcode::G_UMULH:
1153 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1155 case TargetOpcode::G_SADDSAT:
1156 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1157 case TargetOpcode::G_UADDSAT:
1158 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1159 case TargetOpcode::G_SSUBSAT:
1160 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1161 case TargetOpcode::G_USUBSAT:
1162 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1164 case TargetOpcode::G_FFREXP:
1165 return selectFrexp(ResVReg, ResType,
I);
1167 case TargetOpcode::G_UADDO:
1168 return selectOverflowArith(ResVReg, ResType,
I,
1169 ResType->
getOpcode() == SPIRV::OpTypeVector
1170 ? SPIRV::OpIAddCarryV
1171 : SPIRV::OpIAddCarryS);
1172 case TargetOpcode::G_USUBO:
1173 return selectOverflowArith(ResVReg, ResType,
I,
1174 ResType->
getOpcode() == SPIRV::OpTypeVector
1175 ? SPIRV::OpISubBorrowV
1176 : SPIRV::OpISubBorrowS);
1177 case TargetOpcode::G_UMULO:
1178 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1179 case TargetOpcode::G_SMULO:
1180 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1182 case TargetOpcode::G_SEXT:
1183 return selectExt(ResVReg, ResType,
I,
true);
1184 case TargetOpcode::G_ANYEXT:
1185 case TargetOpcode::G_ZEXT:
1186 return selectExt(ResVReg, ResType,
I,
false);
1187 case TargetOpcode::G_TRUNC:
1188 return selectTrunc(ResVReg, ResType,
I);
1189 case TargetOpcode::G_FPTRUNC:
1190 case TargetOpcode::G_FPEXT:
1191 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1193 case TargetOpcode::G_PTRTOINT:
1194 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1195 case TargetOpcode::G_INTTOPTR:
1196 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1197 case TargetOpcode::G_BITCAST:
1198 return selectBitcast(ResVReg, ResType,
I);
1199 case TargetOpcode::G_ADDRSPACE_CAST:
1200 return selectAddrSpaceCast(ResVReg, ResType,
I);
1201 case TargetOpcode::G_PTR_ADD: {
1203 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1207 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1208 (*II).getOpcode() == TargetOpcode::COPY ||
1209 (*II).getOpcode() == SPIRV::OpVariable) &&
1210 getImm(
I.getOperand(2), MRI));
1212 bool IsGVInit =
false;
1216 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1217 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1218 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1219 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1229 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1242 "incompatible result and operand types in a bitcast");
1244 MachineInstrBuilder MIB =
1245 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1252 : SPIRV::OpInBoundsPtrAccessChain))
1256 .
addUse(
I.getOperand(2).getReg())
1259 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1263 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1265 .
addUse(
I.getOperand(2).getReg())
1274 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1277 .
addImm(
static_cast<uint32_t
>(
1278 SPIRV::Opcode::InBoundsPtrAccessChain))
1281 .
addUse(
I.getOperand(2).getReg());
1286 case TargetOpcode::G_ATOMICRMW_OR:
1287 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1288 case TargetOpcode::G_ATOMICRMW_ADD:
1289 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1290 case TargetOpcode::G_ATOMICRMW_AND:
1291 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1292 case TargetOpcode::G_ATOMICRMW_MAX:
1293 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1294 case TargetOpcode::G_ATOMICRMW_MIN:
1295 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1296 case TargetOpcode::G_ATOMICRMW_SUB:
1297 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1298 case TargetOpcode::G_ATOMICRMW_XOR:
1299 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1300 case TargetOpcode::G_ATOMICRMW_UMAX:
1301 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1302 case TargetOpcode::G_ATOMICRMW_UMIN:
1303 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1304 case TargetOpcode::G_ATOMICRMW_XCHG:
1305 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1306 case TargetOpcode::G_ATOMIC_CMPXCHG:
1307 return selectAtomicCmpXchg(ResVReg, ResType,
I);
1309 case TargetOpcode::G_ATOMICRMW_FADD:
1310 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1311 case TargetOpcode::G_ATOMICRMW_FSUB:
1313 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1314 ResType->
getOpcode() == SPIRV::OpTypeVector
1316 : SPIRV::OpFNegate);
1317 case TargetOpcode::G_ATOMICRMW_FMIN:
1318 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1319 case TargetOpcode::G_ATOMICRMW_FMAX:
1320 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1322 case TargetOpcode::G_FENCE:
1323 return selectFence(
I);
1325 case TargetOpcode::G_STACKSAVE:
1326 return selectStackSave(ResVReg, ResType,
I);
1327 case TargetOpcode::G_STACKRESTORE:
1328 return selectStackRestore(
I);
1330 case TargetOpcode::G_UNMERGE_VALUES:
1336 case TargetOpcode::G_TRAP:
1337 case TargetOpcode::G_UBSANTRAP:
1338 case TargetOpcode::DBG_LABEL:
1340 case TargetOpcode::G_DEBUGTRAP:
1341 return selectDebugTrap(ResVReg, ResType,
I);
1348bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1349 SPIRVTypeInst ResType,
1350 MachineInstr &
I)
const {
1351 unsigned Opcode = SPIRV::OpNop;
1358bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1359 SPIRVTypeInst ResType,
1361 GL::GLSLExtInst GLInst,
1362 bool setMIFlags,
bool useMISrc,
1365 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1366 std::string DiagMsg;
1367 raw_string_ostream OS(DiagMsg);
1368 I.print(OS,
true,
false,
false,
false);
1369 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1372 return selectExtInst(ResVReg, ResType,
I,
1373 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1374 setMIFlags, useMISrc, SrcRegs);
1377bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1378 SPIRVTypeInst ResType,
1380 CL::OpenCLExtInst CLInst,
1381 bool setMIFlags,
bool useMISrc,
1383 return selectExtInst(ResVReg, ResType,
I,
1384 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1385 setMIFlags, useMISrc, SrcRegs);
1388bool SPIRVInstructionSelector::selectExtInst(
1389 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1390 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1392 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1393 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1394 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1398bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1399 SPIRVTypeInst ResType,
1402 bool setMIFlags,
bool useMISrc,
1405 for (
const auto &[InstructionSet, Opcode] : Insts) {
1409 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1412 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1417 const unsigned NumOps =
I.getNumOperands();
1420 I.getOperand(Index).getType() ==
1421 MachineOperand::MachineOperandType::MO_IntrinsicID)
1424 MIB.
add(
I.getOperand(Index));
1436bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1437 SPIRVTypeInst ResType,
1438 MachineInstr &
I)
const {
1439 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1440 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1441 for (
const auto &Ex : ExtInsts) {
1442 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1443 uint32_t Opcode = Ex.second;
1447 MachineIRBuilder MIRBuilder(
I);
1450 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1455 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1458 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1461 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1464 .
addImm(
static_cast<uint32_t
>(Ex.first))
1466 .
add(
I.getOperand(2))
1470 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1471 .
addDef(
I.getOperand(1).getReg())
1480bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1481 SPIRVTypeInst ResType,
1482 MachineInstr &
I)
const {
1483 Register CosResVReg =
I.getOperand(1).getReg();
1484 unsigned SrcIdx =
I.getNumExplicitDefs();
1489 MachineIRBuilder MIRBuilder(
I);
1491 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1496 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1499 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1501 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1504 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1506 .
add(
I.getOperand(SrcIdx))
1509 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1517 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1520 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1522 .
add(
I.getOperand(SrcIdx))
1524 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1527 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1529 .
add(
I.getOperand(SrcIdx))
1536bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1537 SPIRVTypeInst ResType,
1539 std::vector<Register> Srcs,
1540 unsigned Opcode)
const {
1541 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1551bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1552 SPIRVTypeInst ResType,
1554 unsigned Opcode)
const {
1556 Register SrcReg =
I.getOperand(1).getReg();
1561 unsigned DefOpCode = DefIt->getOpcode();
1562 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1565 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1566 DefOpCode = VRD->getOpcode();
1568 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1569 DefOpCode == TargetOpcode::G_CONSTANT ||
1570 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1576 uint32_t SpecOpcode = 0;
1578 case SPIRV::OpConvertPtrToU:
1579 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1581 case SPIRV::OpConvertUToPtr:
1582 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1587 TII.get(SPIRV::OpSpecConstantOp))
1597 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1601bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1602 SPIRVTypeInst ResType,
1603 MachineInstr &
I)
const {
1604 Register OpReg =
I.getOperand(1).getReg();
1605 SPIRVTypeInst OpType =
1609 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1619 if (
MemOp->isVolatile())
1620 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1621 if (
MemOp->isNonTemporal())
1622 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1624 if (!ST->isShader() &&
MemOp->getAlign().value())
1625 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1629 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1630 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1634 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1636 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1640 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1644 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1646 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1658 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1660 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1662 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1666bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1667 SPIRVTypeInst ResType,
1668 MachineInstr &
I)
const {
1670 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1675 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1676 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1678 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1682 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1686 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1687 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1688 I.getDebugLoc(),
I);
1692 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1696 if (!
I.getNumMemOperands()) {
1697 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1699 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1702 MachineIRBuilder MIRBuilder(
I);
1709bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1711 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1712 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1717 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1718 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1723 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1727 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1728 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1729 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1730 TII.get(SPIRV::OpImageWrite))
1736 if (sampledTypeIsSignedInteger(LLVMHandleType))
1739 BMI.constrainAllUses(
TII,
TRI, RBI);
1745 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpStore))
1748 if (!
I.getNumMemOperands()) {
1749 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1751 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1754 MachineIRBuilder MIRBuilder(
I);
1761bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
1762 SPIRVTypeInst ResType,
1763 MachineInstr &
I)
const {
1764 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
1772 const Register PtrsReg =
I.getOperand(2).getReg();
1773 const uint32_t Alignment =
I.getOperand(3).getImm();
1774 const Register MaskReg =
I.getOperand(4).getReg();
1775 const Register PassthruReg =
I.getOperand(5).getReg();
1776 const Register AlignmentReg = buildI32Constant(Alignment,
I);
1780 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
1791bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
1792 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
1799 const Register ValuesReg =
I.getOperand(1).getReg();
1800 const Register PtrsReg =
I.getOperand(2).getReg();
1801 const uint32_t Alignment =
I.getOperand(3).getImm();
1802 const Register MaskReg =
I.getOperand(4).getReg();
1803 const Register AlignmentReg = buildI32Constant(Alignment,
I);
1807 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
1816bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
1817 const Twine &Msg)
const {
1818 const Function &
F =
I.getMF()->getFunction();
1819 F.getContext().diagnose(
1820 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
1824bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
1825 SPIRVTypeInst ResType,
1826 MachineInstr &
I)
const {
1827 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1829 "llvm.stacksave intrinsic: this instruction requires the following "
1830 "SPIR-V extension: SPV_INTEL_variable_length_array",
1833 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
1840bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
1841 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1843 "llvm.stackrestore intrinsic: this instruction requires the following "
1844 "SPIR-V extension: SPV_INTEL_variable_length_array",
1846 if (!
I.getOperand(0).isReg())
1849 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
1850 .
addUse(
I.getOperand(0).getReg())
1856SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
1857 MachineIRBuilder MIRBuilder(
I);
1858 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1865 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
1869 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1870 Type *ArrTy = ArrayType::get(ValTy, Num);
1872 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
1875 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
1882 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
1885 .
addImm(SPIRV::StorageClass::UniformConstant)
1896bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
1899 Register DstReg =
I.getOperand(0).getReg();
1904 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
1909 "Unable to determine pointee type size for OpCopyMemory");
1910 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
1911 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
1913 "OpCopyMemory requires the size to match the pointee type size");
1914 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
1917 if (
I.getNumMemOperands()) {
1918 MachineIRBuilder MIRBuilder(
I);
1925bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
1928 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
1929 .
addUse(
I.getOperand(0).getReg())
1931 .
addUse(
I.getOperand(2).getReg());
1932 if (
I.getNumMemOperands()) {
1933 MachineIRBuilder MIRBuilder(
I);
1940bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
1941 MachineInstr &
I)
const {
1942 Register SrcReg =
I.getOperand(1).getReg();
1943 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
1944 Register VarReg = getOrCreateMemSetGlobal(
I);
1947 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1949 ValTy,
I, SPIRV::StorageClass::UniformConstant);
1951 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
1955 if (!selectCopyMemory(
I, SrcReg))
1958 if (!selectCopyMemorySized(
I, SrcReg))
1961 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
1962 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
1967bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
1968 SPIRVTypeInst ResType,
1971 unsigned NegateOpcode)
const {
1973 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1976 Register ScopeReg = buildI32Constant(Scope,
I);
1978 Register Ptr =
I.getOperand(1).getReg();
1984 Register MemSemReg = buildI32Constant(MemSem ,
I);
1986 Register ValueReg =
I.getOperand(2).getReg();
1987 if (NegateOpcode != 0) {
1990 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
1995 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2006bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2007 unsigned ArgI =
I.getNumOperands() - 1;
2009 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2010 SPIRVTypeInst SrcType =
2012 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2014 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2016 SPIRVTypeInst ScalarType =
2019 unsigned CurrentIndex = 0;
2020 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2021 Register ResVReg =
I.getOperand(i).getReg();
2024 LLT ResLLT = MRI->
getType(ResVReg);
2030 ResType = ScalarType;
2036 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2039 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2045 for (
unsigned j = 0;
j < NumElements; ++
j) {
2046 MIB.
addImm(CurrentIndex + j);
2048 CurrentIndex += NumElements;
2052 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2064bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2067 Register MemSemReg = buildI32Constant(MemSem,
I);
2069 uint32_t
Scope =
static_cast<uint32_t
>(
2071 Register ScopeReg = buildI32Constant(Scope,
I);
2073 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2080bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2081 SPIRVTypeInst ResType,
2083 unsigned Opcode)
const {
2084 Type *ResTy =
nullptr;
2088 "Not enough info to select the arithmetic with overflow instruction");
2091 "with overflow instruction");
2097 MachineIRBuilder MIRBuilder(
I);
2099 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2100 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2106 Register ZeroReg = buildZerosVal(ResType,
I);
2111 if (ResName.
size() > 0)
2116 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2119 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2120 MIB.
addUse(
I.getOperand(i).getReg());
2125 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2126 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2128 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2129 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2136 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2137 .
addDef(
I.getOperand(1).getReg())
2145bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2146 SPIRVTypeInst ResType,
2147 MachineInstr &
I)
const {
2151 Register Ptr =
I.getOperand(2).getReg();
2154 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2157 ScopeReg = buildI32Constant(Scope,
I);
2159 unsigned ScSem =
static_cast<uint32_t
>(
2162 unsigned MemSemEq =
static_cast<uint32_t
>(
getMemSemantics(AO)) | ScSem;
2163 Register MemSemEqReg = buildI32Constant(MemSemEq,
I);
2165 unsigned MemSemNeq =
static_cast<uint32_t
>(
getMemSemantics(FO)) | ScSem;
2166 if (MemSemEq == MemSemNeq)
2167 MemSemNeqReg = MemSemEqReg;
2169 MemSemNeqReg = buildI32Constant(MemSemEq,
I);
2172 ScopeReg =
I.getOperand(5).getReg();
2173 MemSemEqReg =
I.getOperand(6).getReg();
2174 MemSemNeqReg =
I.getOperand(7).getReg();
2178 Register Val =
I.getOperand(4).getReg();
2182 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2201 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2208 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2220 case SPIRV::StorageClass::DeviceOnlyINTEL:
2221 case SPIRV::StorageClass::HostOnlyINTEL:
2230 bool IsGRef =
false;
2231 bool IsAllowedRefs =
2233 unsigned Opcode = It.getOpcode();
2234 if (Opcode == SPIRV::OpConstantComposite ||
2235 Opcode == SPIRV::OpSpecConstantComposite ||
2236 Opcode == SPIRV::OpVariable ||
2237 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2238 return IsGRef = true;
2239 return Opcode == SPIRV::OpName;
2241 return IsAllowedRefs && IsGRef;
2244Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2245 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2247 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2251SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2253 uint32_t Opcode)
const {
2254 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2255 TII.get(SPIRV::OpSpecConstantOp))
2263SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2264 SPIRVTypeInst SrcPtrTy)
const {
2265 SPIRVTypeInst GenericPtrTy =
2269 SPIRV::StorageClass::Generic),
2271 MachineFunction *MF =
I.getParent()->getParent();
2273 MachineInstrBuilder MIB = buildSpecConstantOp(
2275 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2285bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2286 SPIRVTypeInst ResType,
2287 MachineInstr &
I)
const {
2291 Register SrcPtr =
I.getOperand(1).getReg();
2295 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2296 ResType->
getOpcode() != SPIRV::OpTypePointer)
2297 return BuildCOPY(ResVReg, SrcPtr,
I);
2307 unsigned SpecOpcode =
2309 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2312 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2319 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2321 .constrainAllUses(
TII,
TRI, RBI);
2323 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2325 buildSpecConstantOp(
2327 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2328 .constrainAllUses(
TII,
TRI, RBI);
2335 return BuildCOPY(ResVReg, SrcPtr,
I);
2337 if ((SrcSC == SPIRV::StorageClass::Function &&
2338 DstSC == SPIRV::StorageClass::Private) ||
2339 (DstSC == SPIRV::StorageClass::Function &&
2340 SrcSC == SPIRV::StorageClass::Private))
2341 return BuildCOPY(ResVReg, SrcPtr,
I);
2345 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2348 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2351 SPIRVTypeInst GenericPtrTy =
2370 return selectUnOp(ResVReg, ResType,
I,
2371 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2373 return selectUnOp(ResVReg, ResType,
I,
2374 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2376 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2378 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2388 return SPIRV::OpFOrdEqual;
2390 return SPIRV::OpFOrdGreaterThanEqual;
2392 return SPIRV::OpFOrdGreaterThan;
2394 return SPIRV::OpFOrdLessThanEqual;
2396 return SPIRV::OpFOrdLessThan;
2398 return SPIRV::OpFOrdNotEqual;
2400 return SPIRV::OpOrdered;
2402 return SPIRV::OpFUnordEqual;
2404 return SPIRV::OpFUnordGreaterThanEqual;
2406 return SPIRV::OpFUnordGreaterThan;
2408 return SPIRV::OpFUnordLessThanEqual;
2410 return SPIRV::OpFUnordLessThan;
2412 return SPIRV::OpFUnordNotEqual;
2414 return SPIRV::OpUnordered;
2424 return SPIRV::OpIEqual;
2426 return SPIRV::OpINotEqual;
2428 return SPIRV::OpSGreaterThanEqual;
2430 return SPIRV::OpSGreaterThan;
2432 return SPIRV::OpSLessThanEqual;
2434 return SPIRV::OpSLessThan;
2436 return SPIRV::OpUGreaterThanEqual;
2438 return SPIRV::OpUGreaterThan;
2440 return SPIRV::OpULessThanEqual;
2442 return SPIRV::OpULessThan;
2451 return SPIRV::OpPtrEqual;
2453 return SPIRV::OpPtrNotEqual;
2464 return SPIRV::OpLogicalEqual;
2466 return SPIRV::OpLogicalNotEqual;
2500bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2501 SPIRVTypeInst ResType,
2503 unsigned OpAnyOrAll)
const {
2504 assert(
I.getNumOperands() == 3);
2505 assert(
I.getOperand(2).isReg());
2507 Register InputRegister =
I.getOperand(2).getReg();
2510 assert(InputType &&
"VReg has no type assigned");
2513 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2514 if (IsBoolTy && !IsVectorTy) {
2515 assert(ResVReg ==
I.getOperand(0).getReg());
2516 return BuildCOPY(ResVReg, InputRegister,
I);
2520 unsigned SpirvNotEqualId =
2521 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2523 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2528 IsBoolTy ? InputRegister
2536 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2538 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2555bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2556 SPIRVTypeInst ResType,
2557 MachineInstr &
I)
const {
2558 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2561bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2562 SPIRVTypeInst ResType,
2563 MachineInstr &
I)
const {
2564 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2568bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2569 SPIRVTypeInst ResType,
2570 MachineInstr &
I)
const {
2571 assert(
I.getNumOperands() == 4);
2572 assert(
I.getOperand(2).isReg());
2573 assert(
I.getOperand(3).isReg());
2575 [[maybe_unused]] SPIRVTypeInst VecType =
2580 "dot product requires a vector of at least 2 components");
2582 [[maybe_unused]] SPIRVTypeInst EltType =
2591 .
addUse(
I.getOperand(2).getReg())
2592 .
addUse(
I.getOperand(3).getReg())
2597bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2598 SPIRVTypeInst ResType,
2601 assert(
I.getNumOperands() == 4);
2602 assert(
I.getOperand(2).isReg());
2603 assert(
I.getOperand(3).isReg());
2606 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2610 .
addUse(
I.getOperand(2).getReg())
2611 .
addUse(
I.getOperand(3).getReg())
2618bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2619 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2620 assert(
I.getNumOperands() == 4);
2621 assert(
I.getOperand(2).isReg());
2622 assert(
I.getOperand(3).isReg());
2626 Register Vec0 =
I.getOperand(2).getReg();
2627 Register Vec1 =
I.getOperand(3).getReg();
2631 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2640 "dot product requires a vector of at least 2 components");
2643 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2653 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2664 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2676bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2677 SPIRVTypeInst ResType,
2678 MachineInstr &
I)
const {
2680 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2683 .
addUse(
I.getOperand(2).getReg())
2688bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2689 SPIRVTypeInst ResType,
2690 MachineInstr &
I)
const {
2692 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2695 .
addUse(
I.getOperand(2).getReg())
2700template <
bool Signed>
2701bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2702 SPIRVTypeInst ResType,
2703 MachineInstr &
I)
const {
2704 assert(
I.getNumOperands() == 5);
2705 assert(
I.getOperand(2).isReg());
2706 assert(
I.getOperand(3).isReg());
2707 assert(
I.getOperand(4).isReg());
2710 Register Acc =
I.getOperand(2).getReg();
2714 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2716 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
2721 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
2724 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2736template <
bool Signed>
2737bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
2738 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2739 assert(
I.getNumOperands() == 5);
2740 assert(
I.getOperand(2).isReg());
2741 assert(
I.getOperand(3).isReg());
2742 assert(
I.getOperand(4).isReg());
2745 Register Acc =
I.getOperand(2).getReg();
2751 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
2755 for (
unsigned i = 0; i < 4; i++) {
2778 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
2798 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2813bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
2814 SPIRVTypeInst ResType,
2815 MachineInstr &
I)
const {
2816 assert(
I.getNumOperands() == 3);
2817 assert(
I.getOperand(2).isReg());
2819 Register VZero = buildZerosValF(ResType,
I);
2820 Register VOne = buildOnesValF(ResType,
I);
2822 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
2825 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2827 .
addUse(
I.getOperand(2).getReg())
2834bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
2835 SPIRVTypeInst ResType,
2836 MachineInstr &
I)
const {
2837 assert(
I.getNumOperands() == 3);
2838 assert(
I.getOperand(2).isReg());
2840 Register InputRegister =
I.getOperand(2).getReg();
2842 auto &
DL =
I.getDebugLoc();
2852 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
2854 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
2862 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2867 if (NeedsConversion) {
2868 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
2879bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
2880 SPIRVTypeInst ResType,
2882 unsigned Opcode)
const {
2886 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2892 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
2893 BMI.addUse(
I.getOperand(J).getReg());
2900bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
2903 bool WithGroupSync)
const {
2905 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
2907 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
2909 assert(((Scope != SPIRV::Scope::Workgroup) ||
2910 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
2911 "Workgroup Scope must set WorkGroupMemory semantic "
2912 "in Barrier instruction");
2914 assert(((Scope != SPIRV::Scope::Device) ||
2915 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
2916 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
2917 "Device Scope must set UniformMemory and ImageMemory semantic "
2918 "in Barrier instruction");
2920 Register MemSemReg = buildI32Constant(MemSem,
I);
2921 Register ScopeReg = buildI32Constant(Scope,
I);
2927 if (WithGroupSync) {
2928 MI.addUse(ScopeReg);
2931 MI.addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
2935bool SPIRVInstructionSelector::selectWaveActiveCountBits(
2936 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2941 if (!selectWaveOpInst(BallotReg, BallotType,
I,
2942 SPIRV::OpGroupNonUniformBallot))
2947 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2952 .
addImm(SPIRV::GroupOperation::Reduce)
2961 if (
Type->getOpcode() != SPIRV::OpTypeVector)
2965 return Type->getOperand(2).getImm();
2968bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
2969 SPIRVTypeInst ResType,
2970 MachineInstr &
I)
const {
2975 Register InputReg =
I.getOperand(2).getReg();
2980 bool IsVector = NumElems > 1;
2983 SPIRVTypeInst ElemInputType = InputType;
2984 SPIRVTypeInst ElemBoolType = ResType;
2997 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
2998 SPIRV::OpGroupNonUniformAllEqual);
3003 ElementResults.
reserve(NumElems);
3005 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3018 ElemInput = Extracted;
3024 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3035 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3046bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3047 SPIRVTypeInst ResType,
3048 MachineInstr &
I)
const {
3050 assert(
I.getNumOperands() == 3);
3052 auto Op =
I.getOperand(2);
3064 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3086 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3090 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3097bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3098 SPIRVTypeInst ResType,
3100 bool IsUnsigned)
const {
3101 return selectWaveReduce(
3102 ResVReg, ResType,
I, IsUnsigned,
3103 [&](
Register InputRegister,
bool IsUnsigned) {
3104 const bool IsFloatTy =
3106 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3107 : SPIRV::OpGroupNonUniformSMax;
3108 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3112bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3113 SPIRVTypeInst ResType,
3115 bool IsUnsigned)
const {
3116 return selectWaveReduce(
3117 ResVReg, ResType,
I, IsUnsigned,
3118 [&](
Register InputRegister,
bool IsUnsigned) {
3119 const bool IsFloatTy =
3121 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3122 : SPIRV::OpGroupNonUniformSMin;
3123 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3127bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3128 SPIRVTypeInst ResType,
3129 MachineInstr &
I)
const {
3130 return selectWaveReduce(ResVReg, ResType,
I,
false,
3131 [&](
Register InputRegister,
bool IsUnsigned) {
3133 InputRegister, SPIRV::OpTypeFloat);
3134 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3135 : SPIRV::OpGroupNonUniformIAdd;
3139bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3140 SPIRVTypeInst ResType,
3141 MachineInstr &
I)
const {
3142 return selectWaveReduce(ResVReg, ResType,
I,
false,
3143 [&](
Register InputRegister,
bool IsUnsigned) {
3145 InputRegister, SPIRV::OpTypeFloat);
3146 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3147 : SPIRV::OpGroupNonUniformIMul;
3151template <
typename PickOpcodeFn>
3152bool SPIRVInstructionSelector::selectWaveReduce(
3153 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3154 PickOpcodeFn &&PickOpcode)
const {
3155 assert(
I.getNumOperands() == 3);
3156 assert(
I.getOperand(2).isReg());
3158 Register InputRegister =
I.getOperand(2).getReg();
3165 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3171 .
addImm(SPIRV::GroupOperation::Reduce)
3172 .
addUse(
I.getOperand(2).getReg())
3177bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3178 SPIRVTypeInst ResType,
3180 unsigned Opcode)
const {
3181 return selectWaveReduce(
3182 ResVReg, ResType,
I,
false,
3183 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3186bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3187 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3188 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3189 [&](
Register InputRegister,
bool IsUnsigned) {
3191 InputRegister, SPIRV::OpTypeFloat);
3193 ? SPIRV::OpGroupNonUniformFAdd
3194 : SPIRV::OpGroupNonUniformIAdd;
3198bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3199 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3200 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3201 [&](
Register InputRegister,
bool IsUnsigned) {
3203 InputRegister, SPIRV::OpTypeFloat);
3205 ? SPIRV::OpGroupNonUniformFMul
3206 : SPIRV::OpGroupNonUniformIMul;
3210template <
typename PickOpcodeFn>
3211bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3212 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3213 PickOpcodeFn &&PickOpcode)
const {
3214 assert(
I.getNumOperands() == 3);
3215 assert(
I.getOperand(2).isReg());
3217 Register InputRegister =
I.getOperand(2).getReg();
3224 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3230 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3231 .
addUse(
I.getOperand(2).getReg())
3236bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3237 SPIRVTypeInst ResType,
3240 assert(
I.getNumOperands() == 3);
3241 assert(
I.getOperand(2).isReg());
3243 Register InputRegister =
I.getOperand(2).getReg();
3249 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3260bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3261 SPIRVTypeInst ResType,
3266 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3271 : SPIRV::OpUConvert;
3275 ShiftOp = SPIRV::OpShiftRightLogicalV;
3280 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3281 TII.get(SPIRV::OpConstantComposite))
3284 for (
unsigned It = 0; It <
N; ++It)
3288 ShiftConst = CompositeReg;
3293 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3298 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3303 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3308 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3311bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
3312 SPIRVTypeInst ResType,
3316 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3324bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3325 SPIRVTypeInst ResType,
3326 MachineInstr &
I)
const {
3327 Register OpReg =
I.getOperand(1).getReg();
3334 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3336 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3341 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3352 unsigned AndOp = SPIRV::OpBitwiseAndS;
3353 unsigned OrOp = SPIRV::OpBitwiseOrS;
3354 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
3355 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
3357 AndOp = SPIRV::OpBitwiseAndV;
3358 OrOp = SPIRV::OpBitwiseOrV;
3359 ShlOp = SPIRV::OpShiftLeftLogicalV;
3360 ShrOp = SPIRV::OpShiftRightLogicalV;
3366 const unsigned Shift) ->
Register {
3374 Register MaskReg = CreateConst(Mask);
3375 Register ShiftReg = CreateConst(Shift);
3382 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
3383 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
3384 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
3385 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
3386 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
3394 uint64_t
Mask = ~0ull;
3395 while ((Shift >>= 1) > 0) {
3402 return BuildCOPY(ResVReg, Result,
I);
3405bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3406 SPIRVTypeInst ResType,
3407 MachineInstr &
I)
const {
3413 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3415 Register OpReg =
I.getOperand(1).getReg();
3416 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
3417 if (
Def->getOpcode() == TargetOpcode::COPY)
3420 switch (
Def->getOpcode()) {
3421 case SPIRV::ASSIGN_TYPE:
3422 if (MachineInstr *AssignToDef =
3424 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3425 Reg =
Def->getOperand(2).getReg();
3428 case SPIRV::OpUndef:
3429 Reg =
Def->getOperand(1).getReg();
3432 unsigned DestOpCode;
3434 DestOpCode = SPIRV::OpConstantNull;
3436 DestOpCode = TargetOpcode::COPY;
3439 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3440 .
addDef(
I.getOperand(0).getReg())
3448bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3449 SPIRVTypeInst ResType,
3450 MachineInstr &
I)
const {
3452 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3454 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3458 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3463 for (
unsigned i =
I.getNumExplicitDefs();
3464 i <
I.getNumExplicitOperands() && IsConst; ++i)
3468 if (!IsConst &&
N < 2)
3470 "There must be at least two constituent operands in a vector");
3473 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3474 TII.get(IsConst ? SPIRV::OpConstantComposite
3475 : SPIRV::OpCompositeConstruct))
3478 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3479 MIB.
addUse(
I.getOperand(i).getReg());
3484bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3485 SPIRVTypeInst ResType,
3486 MachineInstr &
I)
const {
3488 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3490 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3496 if (!
I.getOperand(
OpIdx).isReg())
3503 if (!IsConst &&
N < 2)
3505 "There must be at least two constituent operands in a vector");
3508 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3509 TII.get(IsConst ? SPIRV::OpConstantComposite
3510 : SPIRV::OpCompositeConstruct))
3513 for (
unsigned i = 0; i <
N; ++i)
3519bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3520 SPIRVTypeInst ResType,
3521 MachineInstr &
I)
const {
3526 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3528 Opcode = SPIRV::OpDemoteToHelperInvocation;
3530 Opcode = SPIRV::OpKill;
3532 if (MachineInstr *NextI =
I.getNextNode()) {
3534 NextI->eraseFromParent();
3544bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3545 SPIRVTypeInst ResType,
unsigned CmpOpc,
3546 MachineInstr &
I)
const {
3547 Register Cmp0 =
I.getOperand(2).getReg();
3548 Register Cmp1 =
I.getOperand(3).getReg();
3551 "CMP operands should have the same type");
3552 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3562bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3563 SPIRVTypeInst ResType,
3564 MachineInstr &
I)
const {
3565 auto Pred =
I.getOperand(1).getPredicate();
3568 Register CmpOperand =
I.getOperand(2).getReg();
3575 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3579SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3580 SPIRVTypeInst ResType)
const {
3582 SPIRVTypeInst SpvI32Ty =
3585 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3592 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3595 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3598 .
addImm(APInt(32, Val).getZExtValue());
3600 GR.
add(ConstInt,
MI);
3605bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
3606 SPIRVTypeInst ResType,
3607 MachineInstr &
I)
const {
3609 return selectCmp(ResVReg, ResType, CmpOp,
I);
3612bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
3613 SPIRVTypeInst ResType,
3614 MachineInstr &
I)
const {
3616 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
3623 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
3624 ResType->
getOpcode() != SPIRV::OpTypeFloat)
3627 MachineIRBuilder MIRBuilder(
I);
3629 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
3635 "only float operands supported by GLSL extended math");
3638 MIRBuilder, SpirvScalarType);
3640 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
3641 ? SPIRV::OpVectorTimesScalar
3644 if (!selectOpWithSrcs(ArgReg, ResType,
I,
3645 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
3647 if (!selectExtInst(ResVReg, ResType,
I,
3648 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
3658Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
3659 MachineInstr &
I)
const {
3662 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3667bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
3673 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
3681 if (
Def->getOpcode() == SPIRV::OpConstantNull)
3684 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
3685 Def->getOpcode() == SPIRV::OpConstantI)
3698 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
3699 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
3701 Intrinsic::spv_const_composite)) {
3702 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
3703 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
3704 if (!IsZero(
Def->getOperand(i).getReg()))
3713Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
3714 MachineInstr &
I)
const {
3718 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3723Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
3724 MachineInstr &
I)
const {
3728 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3734 SPIRVTypeInst ResType,
3735 MachineInstr &
I)
const {
3739 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3744bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
3745 SPIRVTypeInst ResType,
3746 MachineInstr &
I)
const {
3747 Register SelectFirstArg =
I.getOperand(2).getReg();
3748 Register SelectSecondArg =
I.getOperand(3).getReg();
3757 SPIRV::OpTypeVector;
3764 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
3765 }
else if (IsPtrTy) {
3766 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
3768 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
3772 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
3773 }
else if (IsPtrTy) {
3774 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
3776 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3779 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3782 .
addUse(
I.getOperand(1).getReg())
3791bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
3792 SPIRVTypeInst ResType,
3794 MachineInstr &InsertAt,
3795 bool IsSigned)
const {
3797 Register ZeroReg = buildZerosVal(ResType, InsertAt);
3798 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
3799 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
3801 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3813bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
3814 SPIRVTypeInst ResType,
3815 MachineInstr &
I,
bool IsSigned,
3816 unsigned Opcode)
const {
3817 Register SrcReg =
I.getOperand(1).getReg();
3823 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
3828 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
3830 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
3833bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
3834 SPIRVTypeInst ResType, MachineInstr &
I,
3835 bool IsSigned)
const {
3836 Register SrcReg =
I.getOperand(1).getReg();
3838 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
3842 if (ResType == SrcType)
3843 return BuildCOPY(ResVReg, SrcReg,
I);
3845 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3846 return selectUnOp(ResVReg, ResType,
I, Opcode);
3849bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
3850 SPIRVTypeInst ResType,
3852 bool IsSigned)
const {
3853 MachineIRBuilder MIRBuilder(
I);
3854 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
3869 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
3872 .
addUse(
I.getOperand(1).getReg())
3873 .
addUse(
I.getOperand(2).getReg())
3879 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
3882 .
addUse(
I.getOperand(1).getReg())
3883 .
addUse(
I.getOperand(2).getReg())
3891 unsigned SelectOpcode =
3892 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
3897 .
addUse(buildOnesVal(
true, ResType,
I))
3898 .
addUse(buildZerosVal(ResType,
I))
3905 .
addUse(buildOnesVal(
false, ResType,
I))
3910bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
3913 SPIRVTypeInst IntTy,
3914 SPIRVTypeInst BoolTy)
const {
3917 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
3918 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
3920 Register One = buildOnesVal(
false, IntTy,
I);
3928 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
3937bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
3938 SPIRVTypeInst ResType,
3939 MachineInstr &
I)
const {
3940 Register IntReg =
I.getOperand(1).getReg();
3943 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
3944 if (ArgType == ResType)
3945 return BuildCOPY(ResVReg, IntReg,
I);
3947 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3948 return selectUnOp(ResVReg, ResType,
I, Opcode);
3951bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
3952 SPIRVTypeInst ResType,
3953 MachineInstr &
I)
const {
3954 unsigned Opcode =
I.getOpcode();
3955 unsigned TpOpcode = ResType->
getOpcode();
3957 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
3958 assert(Opcode == TargetOpcode::G_CONSTANT &&
3959 I.getOperand(1).getCImm()->isZero());
3960 MachineBasicBlock &DepMBB =
I.getMF()->front();
3963 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
3970 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
3973bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
3974 SPIRVTypeInst ResType,
3975 MachineInstr &
I)
const {
3976 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3983bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
3984 SPIRVTypeInst ResType,
3985 MachineInstr &
I)
const {
3987 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
3991 .
addUse(
I.getOperand(3).getReg())
3993 .
addUse(
I.getOperand(2).getReg());
3994 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4000bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4001 SPIRVTypeInst ResType,
4002 MachineInstr &
I)
const {
4003 Type *MaybeResTy =
nullptr;
4008 "Expected aggregate type for extractv instruction");
4010 SPIRV::AccessQualifier::ReadWrite,
false);
4014 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4017 .
addUse(
I.getOperand(2).getReg());
4018 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4024bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4025 SPIRVTypeInst ResType,
4026 MachineInstr &
I)
const {
4027 if (
getImm(
I.getOperand(4), MRI))
4028 return selectInsertVal(ResVReg, ResType,
I);
4030 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4033 .
addUse(
I.getOperand(2).getReg())
4034 .
addUse(
I.getOperand(3).getReg())
4035 .
addUse(
I.getOperand(4).getReg())
4040bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4041 SPIRVTypeInst ResType,
4042 MachineInstr &
I)
const {
4043 if (
getImm(
I.getOperand(3), MRI))
4044 return selectExtractVal(ResVReg, ResType,
I);
4046 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4049 .
addUse(
I.getOperand(2).getReg())
4050 .
addUse(
I.getOperand(3).getReg())
4055bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4056 SPIRVTypeInst ResType,
4057 MachineInstr &
I)
const {
4058 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4064 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4065 : SPIRV::OpAccessChain)
4066 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4067 :
SPIRV::OpPtrAccessChain);
4069 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4073 .
addUse(
I.getOperand(3).getReg());
4075 (Opcode == SPIRV::OpPtrAccessChain ||
4076 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4077 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4078 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4081 const unsigned StartingIndex =
4082 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4085 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4086 Res.addUse(
I.getOperand(i).getReg());
4087 Res.constrainAllUses(
TII,
TRI, RBI);
4092bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4094 unsigned Lim =
I.getNumExplicitOperands();
4095 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4096 Register OpReg =
I.getOperand(i).getReg();
4097 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4099 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4100 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4101 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4108 MachineFunction *MF =
I.getMF();
4120 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4121 TII.get(SPIRV::OpSpecConstantOp))
4124 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4126 GR.
add(OpDefine, MIB);
4132bool SPIRVInstructionSelector::selectDerivativeInst(
4133 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4134 const unsigned DPdOpCode)
const {
4137 errorIfInstrOutsideShader(
I);
4142 Register SrcReg =
I.getOperand(2).getReg();
4147 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4150 .
addUse(
I.getOperand(2).getReg());
4152 MachineIRBuilder MIRBuilder(
I);
4155 if (componentCount != 1)
4159 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4163 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4168 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4173 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4181bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4182 SPIRVTypeInst ResType,
4183 MachineInstr &
I)
const {
4187 case Intrinsic::spv_load:
4188 return selectLoad(ResVReg, ResType,
I);
4189 case Intrinsic::spv_store:
4190 return selectStore(
I);
4191 case Intrinsic::spv_extractv:
4192 return selectExtractVal(ResVReg, ResType,
I);
4193 case Intrinsic::spv_insertv:
4194 return selectInsertVal(ResVReg, ResType,
I);
4195 case Intrinsic::spv_extractelt:
4196 return selectExtractElt(ResVReg, ResType,
I);
4197 case Intrinsic::spv_insertelt:
4198 return selectInsertElt(ResVReg, ResType,
I);
4199 case Intrinsic::spv_gep:
4200 return selectGEP(ResVReg, ResType,
I);
4201 case Intrinsic::spv_bitcast: {
4202 Register OpReg =
I.getOperand(2).getReg();
4203 SPIRVTypeInst OpType =
4207 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4209 case Intrinsic::spv_unref_global:
4210 case Intrinsic::spv_init_global: {
4211 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4216 Register GVarVReg =
MI->getOperand(0).getReg();
4217 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4222 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4224 MI->eraseFromParent();
4228 case Intrinsic::spv_undef: {
4229 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4235 case Intrinsic::spv_named_boolean_spec_constant: {
4236 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4237 : SPIRV::OpSpecConstantFalse;
4239 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4240 .
addDef(
I.getOperand(0).getReg())
4243 unsigned SpecId =
I.getOperand(2).getImm();
4245 SPIRV::Decoration::SpecId, {SpecId});
4249 case Intrinsic::spv_const_composite: {
4251 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4257 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4259 std::function<bool(
Register)> HasSpecConstOperand =
4269 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
4270 J < Def->getNumExplicitOperands(); ++J) {
4271 if (
Def->getOperand(J).isReg() &&
4272 HasSpecConstOperand(
Def->getOperand(J).getReg()))
4278 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
4279 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
4280 : SPIRV::OpConstantComposite;
4281 unsigned ContinuedOpc = HasSpecConst
4282 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
4283 : SPIRV::OpConstantCompositeContinuedINTEL;
4284 MachineIRBuilder MIR(
I);
4286 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
4288 for (
auto *Instr : Instructions) {
4289 Instr->setDebugLoc(
I.getDebugLoc());
4294 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4301 case Intrinsic::spv_assign_name: {
4302 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4303 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4304 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4305 i <
I.getNumExplicitOperands(); ++i) {
4306 MIB.
addImm(
I.getOperand(i).getImm());
4311 case Intrinsic::spv_switch: {
4312 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4313 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4314 if (
I.getOperand(i).isReg())
4315 MIB.
addReg(
I.getOperand(i).getReg());
4316 else if (
I.getOperand(i).isCImm())
4317 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4318 else if (
I.getOperand(i).isMBB())
4319 MIB.
addMBB(
I.getOperand(i).getMBB());
4326 case Intrinsic::spv_loop_merge: {
4327 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4328 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4329 if (
I.getOperand(i).isMBB())
4330 MIB.
addMBB(
I.getOperand(i).getMBB());
4337 case Intrinsic::spv_loop_control_intel: {
4339 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4340 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4345 case Intrinsic::spv_selection_merge: {
4347 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4348 assert(
I.getOperand(1).isMBB() &&
4349 "operand 1 to spv_selection_merge must be a basic block");
4350 MIB.
addMBB(
I.getOperand(1).getMBB());
4351 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4355 case Intrinsic::spv_cmpxchg:
4356 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4357 case Intrinsic::spv_unreachable:
4358 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4361 case Intrinsic::spv_alloca:
4362 return selectFrameIndex(ResVReg, ResType,
I);
4363 case Intrinsic::spv_alloca_array:
4364 return selectAllocaArray(ResVReg, ResType,
I);
4365 case Intrinsic::spv_assume:
4367 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4368 .
addUse(
I.getOperand(1).getReg())
4373 case Intrinsic::spv_expect:
4375 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4378 .
addUse(
I.getOperand(2).getReg())
4379 .
addUse(
I.getOperand(3).getReg())
4384 case Intrinsic::arithmetic_fence:
4385 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4386 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4389 .
addUse(
I.getOperand(2).getReg())
4393 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4395 case Intrinsic::spv_thread_id:
4401 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4403 case Intrinsic::spv_thread_id_in_group:
4409 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4411 case Intrinsic::spv_group_id:
4417 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4419 case Intrinsic::spv_flattened_thread_id_in_group:
4426 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4428 case Intrinsic::spv_workgroup_size:
4429 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4431 case Intrinsic::spv_global_size:
4432 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4434 case Intrinsic::spv_global_offset:
4435 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4437 case Intrinsic::spv_num_workgroups:
4438 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4440 case Intrinsic::spv_subgroup_size:
4441 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4443 case Intrinsic::spv_num_subgroups:
4444 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4446 case Intrinsic::spv_subgroup_id:
4447 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4448 case Intrinsic::spv_subgroup_local_invocation_id:
4449 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4450 ResVReg, ResType,
I);
4451 case Intrinsic::spv_subgroup_max_size:
4452 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4454 case Intrinsic::spv_fdot:
4455 return selectFloatDot(ResVReg, ResType,
I);
4456 case Intrinsic::spv_udot:
4457 case Intrinsic::spv_sdot:
4458 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4460 return selectIntegerDot(ResVReg, ResType,
I,
4461 IID == Intrinsic::spv_sdot);
4462 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4463 case Intrinsic::spv_dot4add_i8packed:
4464 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4466 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4467 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4468 case Intrinsic::spv_dot4add_u8packed:
4469 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4471 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4472 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4473 case Intrinsic::spv_all:
4474 return selectAll(ResVReg, ResType,
I);
4475 case Intrinsic::spv_any:
4476 return selectAny(ResVReg, ResType,
I);
4477 case Intrinsic::spv_cross:
4478 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4479 case Intrinsic::spv_distance:
4480 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4481 case Intrinsic::spv_lerp:
4482 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4483 case Intrinsic::spv_length:
4484 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4485 case Intrinsic::spv_degrees:
4486 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4487 case Intrinsic::spv_faceforward:
4488 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4489 case Intrinsic::spv_frac:
4490 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4491 case Intrinsic::spv_isinf:
4492 return selectOpIsInf(ResVReg, ResType,
I);
4493 case Intrinsic::spv_isnan:
4494 return selectOpIsNan(ResVReg, ResType,
I);
4495 case Intrinsic::spv_normalize:
4496 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4497 case Intrinsic::spv_refract:
4498 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4499 case Intrinsic::spv_reflect:
4500 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4501 case Intrinsic::spv_rsqrt:
4502 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4503 case Intrinsic::spv_sign:
4504 return selectSign(ResVReg, ResType,
I);
4505 case Intrinsic::spv_smoothstep:
4506 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4507 case Intrinsic::spv_firstbituhigh:
4508 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4509 case Intrinsic::spv_firstbitshigh:
4510 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4511 case Intrinsic::spv_firstbitlow:
4512 return selectFirstBitLow(ResVReg, ResType,
I);
4513 case Intrinsic::spv_group_memory_barrier:
4514 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4515 SPIRV::MemorySemantics::WorkgroupMemory,
4517 case Intrinsic::spv_group_memory_barrier_with_group_sync:
4518 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4519 SPIRV::MemorySemantics::WorkgroupMemory,
4521 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4522 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4523 SPIRV::StorageClass::StorageClass ResSC =
4527 "Generic storage class");
4528 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4536 case Intrinsic::spv_lifetime_start:
4537 case Intrinsic::spv_lifetime_end: {
4538 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
4539 : SPIRV::OpLifetimeStop;
4540 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
4541 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
4550 case Intrinsic::spv_saturate:
4551 return selectSaturate(ResVReg, ResType,
I);
4552 case Intrinsic::spv_nclamp:
4553 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
4554 case Intrinsic::spv_uclamp:
4555 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
4556 case Intrinsic::spv_sclamp:
4557 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
4558 case Intrinsic::spv_subgroup_prefix_bit_count:
4559 return selectWavePrefixBitCount(ResVReg, ResType,
I);
4560 case Intrinsic::spv_wave_active_countbits:
4561 return selectWaveActiveCountBits(ResVReg, ResType,
I);
4562 case Intrinsic::spv_wave_all_equal:
4563 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
4564 case Intrinsic::spv_wave_all:
4565 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
4566 case Intrinsic::spv_wave_any:
4567 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
4568 case Intrinsic::spv_subgroup_ballot:
4569 return selectWaveOpInst(ResVReg, ResType,
I,
4570 SPIRV::OpGroupNonUniformBallot);
4571 case Intrinsic::spv_wave_is_first_lane:
4572 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
4573 case Intrinsic::spv_wave_reduce_or:
4574 return selectWaveReduceOp(ResVReg, ResType,
I,
4575 SPIRV::OpGroupNonUniformBitwiseOr);
4576 case Intrinsic::spv_wave_reduce_xor:
4577 return selectWaveReduceOp(ResVReg, ResType,
I,
4578 SPIRV::OpGroupNonUniformBitwiseXor);
4579 case Intrinsic::spv_wave_reduce_and:
4580 return selectWaveReduceOp(ResVReg, ResType,
I,
4581 SPIRV::OpGroupNonUniformBitwiseAnd);
4582 case Intrinsic::spv_wave_reduce_umax:
4583 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
4584 case Intrinsic::spv_wave_reduce_max:
4585 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
4586 case Intrinsic::spv_wave_reduce_umin:
4587 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
4588 case Intrinsic::spv_wave_reduce_min:
4589 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
4590 case Intrinsic::spv_wave_reduce_sum:
4591 return selectWaveReduceSum(ResVReg, ResType,
I);
4592 case Intrinsic::spv_wave_product:
4593 return selectWaveReduceProduct(ResVReg, ResType,
I);
4594 case Intrinsic::spv_wave_readlane:
4595 return selectWaveOpInst(ResVReg, ResType,
I,
4596 SPIRV::OpGroupNonUniformShuffle);
4597 case Intrinsic::spv_wave_prefix_sum:
4598 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
4599 case Intrinsic::spv_wave_prefix_product:
4600 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
4601 case Intrinsic::spv_quad_read_across_x: {
4602 return selectQuadSwap(ResVReg, ResType,
I, 0);
4604 case Intrinsic::spv_quad_read_across_y: {
4605 return selectQuadSwap(ResVReg, ResType,
I, 1);
4607 case Intrinsic::spv_step:
4608 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
4609 case Intrinsic::spv_radians:
4610 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
4614 case Intrinsic::instrprof_increment:
4615 case Intrinsic::instrprof_increment_step:
4616 case Intrinsic::instrprof_value_profile:
4619 case Intrinsic::spv_value_md:
4621 case Intrinsic::spv_resource_handlefrombinding: {
4622 return selectHandleFromBinding(ResVReg, ResType,
I);
4624 case Intrinsic::spv_resource_counterhandlefrombinding:
4625 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
4626 case Intrinsic::spv_resource_updatecounter:
4627 return selectUpdateCounter(ResVReg, ResType,
I);
4628 case Intrinsic::spv_resource_store_typedbuffer: {
4629 return selectImageWriteIntrinsic(
I);
4631 case Intrinsic::spv_resource_load_typedbuffer: {
4632 return selectReadImageIntrinsic(ResVReg, ResType,
I);
4634 case Intrinsic::spv_resource_load_level: {
4635 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
4637 case Intrinsic::spv_resource_calculate_lod:
4638 case Intrinsic::spv_resource_calculate_lod_unclamped:
4639 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
4640 case Intrinsic::spv_resource_sample:
4641 case Intrinsic::spv_resource_sample_clamp:
4642 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
4643 case Intrinsic::spv_resource_samplebias:
4644 case Intrinsic::spv_resource_samplebias_clamp:
4645 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
4646 case Intrinsic::spv_resource_samplegrad:
4647 case Intrinsic::spv_resource_samplegrad_clamp:
4648 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
4649 case Intrinsic::spv_resource_samplelevel:
4650 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
4651 case Intrinsic::spv_resource_samplecmp:
4652 case Intrinsic::spv_resource_samplecmp_clamp:
4653 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
4654 case Intrinsic::spv_resource_samplecmplevelzero:
4655 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
4656 case Intrinsic::spv_resource_gather:
4657 case Intrinsic::spv_resource_gather_cmp:
4658 return selectGatherIntrinsic(ResVReg, ResType,
I);
4659 case Intrinsic::spv_resource_getpointer: {
4660 return selectResourceGetPointer(ResVReg, ResType,
I);
4662 case Intrinsic::spv_pushconstant_getpointer: {
4663 return selectPushConstantGetPointer(ResVReg, ResType,
I);
4665 case Intrinsic::spv_discard: {
4666 return selectDiscard(ResVReg, ResType,
I);
4668 case Intrinsic::spv_resource_nonuniformindex: {
4669 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
4671 case Intrinsic::spv_unpackhalf2x16: {
4672 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
4674 case Intrinsic::spv_packhalf2x16: {
4675 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
4677 case Intrinsic::spv_ddx:
4678 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
4679 case Intrinsic::spv_ddy:
4680 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
4681 case Intrinsic::spv_ddx_coarse:
4682 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
4683 case Intrinsic::spv_ddy_coarse:
4684 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
4685 case Intrinsic::spv_ddx_fine:
4686 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
4687 case Intrinsic::spv_ddy_fine:
4688 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
4689 case Intrinsic::spv_fwidth:
4690 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
4691 case Intrinsic::spv_masked_gather:
4692 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4693 return selectMaskedGather(ResVReg, ResType,
I);
4694 return diagnoseUnsupported(
4695 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
4696 case Intrinsic::spv_masked_scatter:
4697 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4698 return selectMaskedScatter(
I);
4699 return diagnoseUnsupported(
4700 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
4702 std::string DiagMsg;
4703 raw_string_ostream OS(DiagMsg);
4705 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
4712bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
4713 SPIRVTypeInst ResType,
4714 MachineInstr &
I)
const {
4717 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
4724bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
4725 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4727 assert(Intr.getIntrinsicID() ==
4728 Intrinsic::spv_resource_counterhandlefrombinding);
4731 Register MainHandleReg = Intr.getOperand(2).getReg();
4733 assert(MainHandleDef->getIntrinsicID() ==
4734 Intrinsic::spv_resource_handlefrombinding);
4738 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
4739 Register IndexReg = MainHandleDef->getOperand(5).getReg();
4740 std::string CounterName =
4745 MachineIRBuilder MIRBuilder(
I);
4747 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
4749 ArraySize, IndexReg, CounterName, MIRBuilder);
4751 return BuildCOPY(ResVReg, CounterVarReg,
I);
4754bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
4755 SPIRVTypeInst ResType,
4756 MachineInstr &
I)
const {
4758 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
4760 Register CounterHandleReg = Intr.getOperand(2).getReg();
4761 Register IncrReg = Intr.getOperand(3).getReg();
4768 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
4769 assert(CounterVarPointeeType &&
4770 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
4771 "Counter variable must be a struct");
4773 SPIRV::StorageClass::StorageBuffer &&
4774 "Counter variable must be in the storage buffer storage class");
4776 "Counter variable must have exactly 1 member in the struct");
4777 const SPIRVTypeInst MemberType =
4780 "Counter variable struct must have a single i32 member");
4784 MachineIRBuilder MIRBuilder(
I);
4786 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
4789 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
4795 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4798 .
addUse(CounterHandleReg)
4805 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
4808 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
4811 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
4820 return BuildCOPY(ResVReg, AtomicRes,
I);
4828 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
4836bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
4837 SPIRVTypeInst ResType,
4838 MachineInstr &
I)
const {
4846 Register ImageReg =
I.getOperand(2).getReg();
4854 Register IdxReg =
I.getOperand(3).getReg();
4856 MachineInstr &Pos =
I;
4858 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
4862bool SPIRVInstructionSelector::generateSampleImage(
4865 DebugLoc Loc, MachineInstr &Pos)
const {
4876 if (!loadHandleBeforePosition(NewSamplerReg,
4882 MachineIRBuilder MIRBuilder(Pos);
4895 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
4896 ImOps.Lod.has_value();
4897 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
4898 : SPIRV::OpImageSampleImplicitLod;
4900 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
4901 : SPIRV::OpImageSampleDrefImplicitLod;
4910 MIB.
addUse(*ImOps.Compare);
4912 uint32_t ImageOperands = 0;
4914 ImageOperands |= SPIRV::ImageOperand::Bias;
4916 ImageOperands |= SPIRV::ImageOperand::Lod;
4917 if (ImOps.GradX && ImOps.GradY)
4918 ImageOperands |= SPIRV::ImageOperand::Grad;
4919 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
4921 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4924 "Non-constant offsets are not supported in sample instructions.");
4928 ImageOperands |= SPIRV::ImageOperand::MinLod;
4930 if (ImageOperands != 0) {
4931 MIB.
addImm(ImageOperands);
4932 if (ImageOperands & SPIRV::ImageOperand::Bias)
4934 if (ImageOperands & SPIRV::ImageOperand::Lod)
4936 if (ImageOperands & SPIRV::ImageOperand::Grad) {
4937 MIB.
addUse(*ImOps.GradX);
4938 MIB.
addUse(*ImOps.GradY);
4941 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4942 MIB.
addUse(*ImOps.Offset);
4943 if (ImageOperands & SPIRV::ImageOperand::MinLod)
4944 MIB.
addUse(*ImOps.MinLod);
4951bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
4952 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4953 Register ImageReg =
I.getOperand(2).getReg();
4954 Register SamplerReg =
I.getOperand(3).getReg();
4955 Register CoordinateReg =
I.getOperand(4).getReg();
4971 if (!loadHandleBeforePosition(
4976 MachineIRBuilder MIRBuilder(
I);
4982 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
4992 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
4999 unsigned ExtractedIndex =
5001 Intrinsic::spv_resource_calculate_lod_unclamped
5005 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5006 TII.get(SPIRV::OpCompositeExtract))
5016bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
5017 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5018 Register ImageReg =
I.getOperand(2).getReg();
5019 Register SamplerReg =
I.getOperand(3).getReg();
5020 Register CoordinateReg =
I.getOperand(4).getReg();
5021 ImageOperands ImOps;
5022 if (
I.getNumOperands() > 5)
5023 ImOps.Offset =
I.getOperand(5).getReg();
5024 if (
I.getNumOperands() > 6)
5025 ImOps.MinLod =
I.getOperand(6).getReg();
5026 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5027 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5030bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5031 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5032 Register ImageReg =
I.getOperand(2).getReg();
5033 Register SamplerReg =
I.getOperand(3).getReg();
5034 Register CoordinateReg =
I.getOperand(4).getReg();
5035 ImageOperands ImOps;
5036 ImOps.Bias =
I.getOperand(5).getReg();
5037 if (
I.getNumOperands() > 6)
5038 ImOps.Offset =
I.getOperand(6).getReg();
5039 if (
I.getNumOperands() > 7)
5040 ImOps.MinLod =
I.getOperand(7).getReg();
5041 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5042 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5045bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
5046 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5047 Register ImageReg =
I.getOperand(2).getReg();
5048 Register SamplerReg =
I.getOperand(3).getReg();
5049 Register CoordinateReg =
I.getOperand(4).getReg();
5050 ImageOperands ImOps;
5051 ImOps.GradX =
I.getOperand(5).getReg();
5052 ImOps.GradY =
I.getOperand(6).getReg();
5053 if (
I.getNumOperands() > 7)
5054 ImOps.Offset =
I.getOperand(7).getReg();
5055 if (
I.getNumOperands() > 8)
5056 ImOps.MinLod =
I.getOperand(8).getReg();
5057 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5058 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5061bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
5062 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5063 Register ImageReg =
I.getOperand(2).getReg();
5064 Register SamplerReg =
I.getOperand(3).getReg();
5065 Register CoordinateReg =
I.getOperand(4).getReg();
5066 ImageOperands ImOps;
5067 ImOps.Lod =
I.getOperand(5).getReg();
5068 if (
I.getNumOperands() > 6)
5069 ImOps.Offset =
I.getOperand(6).getReg();
5070 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5071 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5074bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
5075 SPIRVTypeInst ResType,
5076 MachineInstr &
I)
const {
5077 Register ImageReg =
I.getOperand(2).getReg();
5078 Register SamplerReg =
I.getOperand(3).getReg();
5079 Register CoordinateReg =
I.getOperand(4).getReg();
5080 ImageOperands ImOps;
5081 ImOps.Compare =
I.getOperand(5).getReg();
5082 if (
I.getNumOperands() > 6)
5083 ImOps.Offset =
I.getOperand(6).getReg();
5084 if (
I.getNumOperands() > 7)
5085 ImOps.MinLod =
I.getOperand(7).getReg();
5086 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5087 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5090bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
5091 SPIRVTypeInst ResType,
5092 MachineInstr &
I)
const {
5093 Register ImageReg =
I.getOperand(2).getReg();
5094 Register CoordinateReg =
I.getOperand(3).getReg();
5095 Register LodReg =
I.getOperand(4).getReg();
5097 ImageOperands ImOps;
5099 if (
I.getNumOperands() > 5)
5100 ImOps.Offset =
I.getOperand(5).getReg();
5112 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
5113 I.getDebugLoc(),
I, &ImOps);
5116bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
5117 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5118 Register ImageReg =
I.getOperand(2).getReg();
5119 Register SamplerReg =
I.getOperand(3).getReg();
5120 Register CoordinateReg =
I.getOperand(4).getReg();
5121 ImageOperands ImOps;
5122 ImOps.Compare =
I.getOperand(5).getReg();
5123 if (
I.getNumOperands() > 6)
5124 ImOps.Offset =
I.getOperand(6).getReg();
5127 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5128 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5131bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
5132 SPIRVTypeInst ResType,
5133 MachineInstr &
I)
const {
5134 Register ImageReg =
I.getOperand(2).getReg();
5135 Register SamplerReg =
I.getOperand(3).getReg();
5136 Register CoordinateReg =
I.getOperand(4).getReg();
5139 "ImageReg is not an image type.");
5144 ComponentOrCompareReg =
I.getOperand(5).getReg();
5145 OffsetReg =
I.getOperand(6).getReg();
5148 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
5152 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5153 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
5154 Dim != SPIRV::Dim::DIM_Rect) {
5156 "Gather operations are only supported for 2D, Cube, and Rect images.");
5163 if (!loadHandleBeforePosition(
5168 MachineIRBuilder MIRBuilder(
I);
5169 SPIRVTypeInst SampledImageType =
5174 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5182 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
5184 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
5186 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5191 .
addUse(ComponentOrCompareReg);
5193 uint32_t ImageOperands = 0;
5194 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
5195 if (Dim == SPIRV::Dim::DIM_Cube) {
5197 "Gather operations with offset are not supported for Cube images.");
5201 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5203 ImageOperands |= SPIRV::ImageOperand::Offset;
5207 if (ImageOperands != 0) {
5208 MIB.
addImm(ImageOperands);
5210 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5218bool SPIRVInstructionSelector::generateImageReadOrFetch(
5221 const ImageOperands *ImOps)
const {
5224 "ImageReg is not an image type.");
5226 bool IsSignedInteger =
5231 bool IsFetch = (SampledOp.getImm() == 1);
5233 auto AddOperands = [&](MachineInstrBuilder &MIB) {
5234 uint32_t ImageOperandsMask = 0;
5235 if (IsSignedInteger)
5236 ImageOperandsMask |= 0x1000;
5238 if (IsFetch && ImOps) {
5240 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
5241 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
5243 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
5245 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5249 if (ImageOperandsMask != 0) {
5250 MIB.
addImm(ImageOperandsMask);
5251 if (IsFetch && ImOps) {
5254 if (ImOps->Offset &&
5255 (ImageOperandsMask &
5256 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5257 MIB.
addUse(*ImOps->Offset);
5263 if (ResultSize == 4) {
5266 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5273 BMI.constrainAllUses(
TII,
TRI, RBI);
5277 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5281 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5287 BMI.constrainAllUses(
TII,
TRI, RBI);
5289 if (ResultSize == 1) {
5298 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
5301bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
5302 SPIRVTypeInst ResType,
5303 MachineInstr &
I)
const {
5304 Register ResourcePtr =
I.getOperand(2).getReg();
5306 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
5315 MachineIRBuilder MIRBuilder(
I);
5317 Register IndexReg =
I.getOperand(3).getReg();
5320 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5330bool SPIRVInstructionSelector::selectPushConstantGetPointer(
5331 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5336bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
5337 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5338 Register ObjReg =
I.getOperand(2).getReg();
5339 if (!BuildCOPY(ResVReg, ObjReg,
I))
5349 decorateUsesAsNonUniform(ResVReg);
5353void SPIRVInstructionSelector::decorateUsesAsNonUniform(
5356 while (WorkList.
size() > 0) {
5360 bool IsDecorated =
false;
5362 if (
Use.getOpcode() == SPIRV::OpDecorate &&
5363 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
5369 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
5371 if (ResultReg == CurrentReg)
5379 SPIRV::Decoration::NonUniformEXT, {});
5384bool SPIRVInstructionSelector::extractSubvector(
5386 MachineInstr &InsertionPoint)
const {
5388 [[maybe_unused]] uint64_t InputSize =
5391 assert(InputSize > 1 &&
"The input must be a vector.");
5392 assert(ResultSize > 1 &&
"The result must be a vector.");
5393 assert(ResultSize < InputSize &&
5394 "Cannot extract more element than there are in the input.");
5397 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
5398 for (uint64_t
I = 0;
I < ResultSize;
I++) {
5401 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5410 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
5412 TII.get(SPIRV::OpCompositeConstruct))
5416 for (
Register ComponentReg : ComponentRegisters)
5417 MIB.
addUse(ComponentReg);
5422bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
5423 MachineInstr &
I)
const {
5430 Register ImageReg =
I.getOperand(1).getReg();
5438 Register CoordinateReg =
I.getOperand(2).getReg();
5439 Register DataReg =
I.getOperand(3).getReg();
5442 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
5450Register SPIRVInstructionSelector::buildPointerToResource(
5451 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
5452 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
5453 StringRef Name, MachineIRBuilder MIRBuilder)
const {
5455 if (ArraySize == 1) {
5456 SPIRVTypeInst PtrType =
5459 "SpirvResType did not have an explicit layout.");
5464 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
5465 SPIRVTypeInst VarPointerType =
5468 VarPointerType, Set,
Binding, Name, MIRBuilder);
5470 SPIRVTypeInst ResPointerType =
5483bool SPIRVInstructionSelector::selectFirstBitSet16(
5484 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5485 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
5487 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
5491 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
5494bool SPIRVInstructionSelector::selectFirstBitSet32(
5496 unsigned BitSetOpcode)
const {
5497 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5500 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5507bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
5509 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5516 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
5518 MachineIRBuilder MIRBuilder(
I);
5521 SPIRVTypeInst I64x2Type =
5523 SPIRVTypeInst Vec2ResType =
5526 std::vector<Register> PartialRegs;
5529 unsigned CurrentComponent = 0;
5530 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
5536 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5537 TII.get(SPIRV::OpVectorShuffle))
5542 .
addImm(CurrentComponent)
5543 .
addImm(CurrentComponent + 1);
5550 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType,
I, BitSetResult,
5551 BitSetOpcode, SwapPrimarySide))
5554 PartialRegs.push_back(SubVecBitSetReg);
5558 if (CurrentComponent != ComponentCount) {
5564 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
5565 SPIRV::OpVectorExtractDynamic))
5571 if (!selectFirstBitSet64(FinalElemBitSetReg,
BaseType,
I, FinalElemReg,
5572 BitSetOpcode, SwapPrimarySide))
5575 PartialRegs.push_back(FinalElemBitSetReg);
5580 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
5581 SPIRV::OpCompositeConstruct);
5584bool SPIRVInstructionSelector::selectFirstBitSet64(
5586 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5599 if (ComponentCount > 2) {
5600 return selectFirstBitSet64Overflow(ResVReg, ResType,
I, SrcReg,
5601 BitSetOpcode, SwapPrimarySide);
5605 MachineIRBuilder MIRBuilder(
I);
5607 BaseType, 2 * ComponentCount, MIRBuilder,
false);
5611 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
5617 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
5624 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
5627 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
5628 SPIRV::OpVectorExtractDynamic))
5630 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
5631 SPIRV::OpVectorExtractDynamic))
5635 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5636 TII.get(SPIRV::OpVectorShuffle))
5644 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
5650 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5651 TII.get(SPIRV::OpVectorShuffle))
5659 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
5679 SelectOp = SPIRV::OpSelectSISCond;
5680 AddOp = SPIRV::OpIAddS;
5688 SelectOp = SPIRV::OpSelectVIVCond;
5689 AddOp = SPIRV::OpIAddV;
5695 Register RegSecondaryOffset = Reg0;
5699 if (SwapPrimarySide) {
5700 PrimaryReg = LowReg;
5701 SecondaryReg = HighReg;
5702 RegPrimaryOffset = Reg0;
5703 RegSecondaryOffset = Reg32;
5708 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
5709 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
5714 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
5715 SPIRV::OpINotEqual))
5722 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
5723 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
5728 if (SwapPrimarySide) {
5730 if (!selectOpWithSrcs(RegAdd, ResType,
I,
5731 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
5742 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
5743 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
5748 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
5749 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
5752 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
5756bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
5757 SPIRVTypeInst ResType,
5759 bool IsSigned)
const {
5761 Register OpReg =
I.getOperand(2).getReg();
5764 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5765 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
5769 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5771 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5773 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5777 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
5781bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
5782 SPIRVTypeInst ResType,
5783 MachineInstr &
I)
const {
5785 Register OpReg =
I.getOperand(2).getReg();
5790 unsigned ExtendOpcode = SPIRV::OpUConvert;
5791 unsigned BitSetOpcode = GL::FindILsb;
5795 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5797 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5799 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5806bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
5807 SPIRVTypeInst ResType,
5808 MachineInstr &
I)
const {
5812 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
5815 .
addUse(
I.getOperand(2).getReg())
5818 unsigned Alignment =
I.getOperand(3).getImm();
5824bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
5825 SPIRVTypeInst ResType,
5826 MachineInstr &
I)
const {
5830 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
5833 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
5836 unsigned Alignment =
I.getOperand(2).getImm();
5843bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
5848 const MachineInstr *PrevI =
I.getPrevNode();
5850 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
5854 .
addMBB(
I.getOperand(0).getMBB())
5859 .
addMBB(
I.getOperand(0).getMBB())
5864bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
5875 const MachineInstr *NextI =
I.getNextNode();
5877 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
5883 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
5885 .
addUse(
I.getOperand(0).getReg())
5886 .
addMBB(
I.getOperand(1).getMBB())
5892bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
5893 MachineInstr &
I)
const {
5895 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
5897 const unsigned NumOps =
I.getNumOperands();
5898 for (
unsigned i = 1; i <
NumOps; i += 2) {
5899 MIB.
addUse(
I.getOperand(i + 0).getReg());
5900 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
5906bool SPIRVInstructionSelector::selectGlobalValue(
5907 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
5909 MachineIRBuilder MIRBuilder(
I);
5910 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
5913 std::string GlobalIdent;
5915 unsigned &
ID = UnnamedGlobalIDs[GV];
5917 ID = UnnamedGlobalIDs.
size();
5918 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
5944 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
5951 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
5956 MachineInstrBuilder MIB1 =
5957 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5960 MachineInstrBuilder MIB2 =
5962 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
5966 GR.
add(ConstVal, MIB2);
5974 MachineInstrBuilder MIB3 =
5975 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5978 GR.
add(ConstVal, MIB3);
5982 assert(NewReg != ResVReg);
5983 return BuildCOPY(ResVReg, NewReg,
I);
5993 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
5999 SPIRVTypeInst ResType =
6003 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6008 if (
GlobalVar->isExternallyInitialized() &&
6009 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6010 constexpr unsigned ReadWriteINTEL = 3u;
6013 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6019bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6020 SPIRVTypeInst ResType,
6021 MachineInstr &
I)
const {
6023 return selectExtInst(ResVReg, ResType,
I, CL::log10);
6031 MachineIRBuilder MIRBuilder(
I);
6036 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6039 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6041 .
add(
I.getOperand(1))
6046 ResType->
getOpcode() == SPIRV::OpTypeFloat);
6048 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
6056 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
6057 ? SPIRV::OpVectorTimesScalar
6068bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
6069 SPIRVTypeInst ResType,
6070 MachineInstr &
I)
const {
6073 return selectExtInst(ResVReg, ResType,
I, CL::pown);
6079 Register ExpReg =
I.getOperand(2).getReg();
6081 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
6082 SPIRV::OpConvertSToF))
6084 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
6091bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
6092 SPIRVTypeInst ResType,
6093 MachineInstr &
I)
const {
6109 MachineIRBuilder MIRBuilder(
I);
6112 ResType, MIRBuilder, SPIRV::StorageClass::Function);
6124 MachineBasicBlock &EntryBB =
I.getMF()->front();
6128 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
6131 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
6137 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6140 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
6143 .
add(
I.getOperand(
I.getNumExplicitDefs()))
6147 Register IntegralPartReg =
I.getOperand(1).getReg();
6148 if (IntegralPartReg.
isValid()) {
6150 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6161 assert(
false &&
"GLSL::Modf is deprecated.");
6172bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
6173 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6174 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6175 MachineIRBuilder MIRBuilder(
I);
6176 const SPIRVTypeInst Vec3Ty =
6179 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
6191 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6195 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6201 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6208 assert(
I.getOperand(2).isReg());
6209 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
6213 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6224bool SPIRVInstructionSelector::loadBuiltinInputID(
6225 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6226 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6227 MachineIRBuilder MIRBuilder(
I);
6229 ResType, MIRBuilder, SPIRV::StorageClass::Input);
6244 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6248 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6257SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
6258 MachineInstr &
I)
const {
6259 MachineIRBuilder MIRBuilder(
I);
6260 if (
Type->getOpcode() != SPIRV::OpTypeVector)
6264 if (VectorSize == 4)
6272bool SPIRVInstructionSelector::loadHandleBeforePosition(
6273 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
6274 MachineInstr &Pos)
const {
6277 Intrinsic::spv_resource_handlefrombinding);
6285 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
6286 MachineIRBuilder MIRBuilder(HandleDef);
6287 SPIRVTypeInst VarType = ResType;
6288 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
6290 if (IsStructuredBuffer) {
6295 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
6297 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
6300 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
6301 ArraySize, IndexReg, Name, MIRBuilder);
6305 uint32_t LoadOpcode =
6306 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
6316void SPIRVInstructionSelector::errorIfInstrOutsideShader(
6317 MachineInstr &
I)
const {
6319 std::string DiagMsg;
6320 raw_string_ostream OS(DiagMsg);
6321 I.print(OS,
true,
false,
false,
false);
6322 DiagMsg +=
" is only supported in shaders.\n";
6328InstructionSelector *
6332 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 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)
unsigned getVectorSizeOrOne(SPIRVTypeInst Type)
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 const fltSemantics & IEEEsingle()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
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.
ArrayRef - 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 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
size - 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
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)
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.
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)
MachineBasicBlock::iterator getFirstValidInstructionInsertPoint(MachineBasicBlock &BB)
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)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineInstr &I)
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...