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 selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
387 MachineInstr &
I)
const;
388 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
389 MachineInstr &
I)
const;
390 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
391 MachineInstr &
I)
const;
392 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
393 MachineInstr &
I)
const;
394 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
395 MachineInstr &
I)
const;
396 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
397 SPIRVTypeInst ResType,
398 MachineInstr &
I)
const;
399 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
400 MachineInstr &
I)
const;
401 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
402 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
403 MachineInstr &
I)
const;
404 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
405 MachineInstr &
I)
const;
406 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
407 MachineInstr &
I)
const;
408 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
409 MachineInstr &
I)
const;
410 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
411 MachineInstr &
I)
const;
412 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
414 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
415 MachineInstr &
I)
const;
416 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
417 MachineInstr &
I)
const;
418 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
419 MachineInstr &
I,
const unsigned DPdOpCode)
const;
421 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
422 SPIRVTypeInst ResType =
nullptr)
const;
424 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
425 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
426 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
428 MachineInstr &
I)
const;
429 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
431 bool wrapIntoSpecConstantOp(MachineInstr &
I,
434 Register getUcharPtrTypeReg(MachineInstr &
I,
435 SPIRV::StorageClass::StorageClass SC)
const;
436 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
438 uint32_t Opcode)
const;
439 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
440 SPIRVTypeInst SrcPtrTy)
const;
441 Register buildPointerToResource(SPIRVTypeInst ResType,
442 SPIRV::StorageClass::StorageClass SC,
443 uint32_t Set, uint32_t
Binding,
444 uint32_t ArraySize,
Register IndexReg,
446 MachineIRBuilder MIRBuilder)
const;
447 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
448 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
449 Register &ReadReg, MachineInstr &InsertionPoint)
const;
450 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
453 const ImageOperands *ImOps =
nullptr)
const;
454 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
456 Register CoordinateReg,
const ImageOperands &ImOps,
459 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
460 Register ResVReg, SPIRVTypeInst ResType,
461 MachineInstr &
I)
const;
462 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
463 Register ResVReg, SPIRVTypeInst ResType,
464 MachineInstr &
I)
const;
465 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
466 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
467 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
468 void errorIfInstrOutsideShader(MachineInstr &
I)
const;
471bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
473 if (
TET->getTargetExtName() ==
"spirv.Image") {
476 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
477 return TET->getTypeParameter(0)->isIntegerTy();
481#define GET_GLOBALISEL_IMPL
482#include "SPIRVGenGlobalISel.inc"
483#undef GET_GLOBALISEL_IMPL
489 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
492#include
"SPIRVGenGlobalISel.inc"
495#include
"SPIRVGenGlobalISel.inc"
507 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
511void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
512 if (HasVRegsReset == &MF)
527 for (
const auto &
MBB : MF) {
528 for (
const auto &
MI :
MBB) {
531 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
535 LLT DstType = MRI.
getType(DstReg);
537 LLT SrcType = MRI.
getType(SrcReg);
538 if (DstType != SrcType)
543 if (DstRC != SrcRC && SrcRC)
555 while (!Stack.empty()) {
560 switch (
MI->getOpcode()) {
561 case TargetOpcode::G_INTRINSIC:
562 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
563 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
566 if (IntrID != Intrinsic::spv_const_composite &&
567 IntrID != Intrinsic::spv_undef)
571 case TargetOpcode::G_BUILD_VECTOR:
572 case TargetOpcode::G_SPLAT_VECTOR:
574 i < OpDef->getNumOperands(); i++) {
579 Stack.push_back(OpNestedDef);
582 case TargetOpcode::G_CONSTANT:
583 case TargetOpcode::G_FCONSTANT:
584 case TargetOpcode::G_IMPLICIT_DEF:
585 case SPIRV::OpConstantTrue:
586 case SPIRV::OpConstantFalse:
587 case SPIRV::OpConstantI:
588 case SPIRV::OpConstantF:
589 case SPIRV::OpConstantComposite:
590 case SPIRV::OpConstantCompositeContinuedINTEL:
591 case SPIRV::OpConstantSampler:
592 case SPIRV::OpConstantNull:
594 case SPIRV::OpConstantFunctionPointerINTEL:
621 case Intrinsic::spv_all:
622 case Intrinsic::spv_alloca:
623 case Intrinsic::spv_any:
624 case Intrinsic::spv_bitcast:
625 case Intrinsic::spv_const_composite:
626 case Intrinsic::spv_cross:
627 case Intrinsic::spv_degrees:
628 case Intrinsic::spv_distance:
629 case Intrinsic::spv_extractelt:
630 case Intrinsic::spv_extractv:
631 case Intrinsic::spv_faceforward:
632 case Intrinsic::spv_fdot:
633 case Intrinsic::spv_firstbitlow:
634 case Intrinsic::spv_firstbitshigh:
635 case Intrinsic::spv_firstbituhigh:
636 case Intrinsic::spv_frac:
637 case Intrinsic::spv_gep:
638 case Intrinsic::spv_global_offset:
639 case Intrinsic::spv_global_size:
640 case Intrinsic::spv_group_id:
641 case Intrinsic::spv_insertelt:
642 case Intrinsic::spv_insertv:
643 case Intrinsic::spv_isinf:
644 case Intrinsic::spv_isnan:
645 case Intrinsic::spv_lerp:
646 case Intrinsic::spv_length:
647 case Intrinsic::spv_normalize:
648 case Intrinsic::spv_num_subgroups:
649 case Intrinsic::spv_num_workgroups:
650 case Intrinsic::spv_ptrcast:
651 case Intrinsic::spv_radians:
652 case Intrinsic::spv_reflect:
653 case Intrinsic::spv_refract:
654 case Intrinsic::spv_resource_getpointer:
655 case Intrinsic::spv_resource_handlefrombinding:
656 case Intrinsic::spv_resource_handlefromimplicitbinding:
657 case Intrinsic::spv_resource_nonuniformindex:
658 case Intrinsic::spv_resource_sample:
659 case Intrinsic::spv_rsqrt:
660 case Intrinsic::spv_saturate:
661 case Intrinsic::spv_sdot:
662 case Intrinsic::spv_sign:
663 case Intrinsic::spv_smoothstep:
664 case Intrinsic::spv_step:
665 case Intrinsic::spv_subgroup_id:
666 case Intrinsic::spv_subgroup_local_invocation_id:
667 case Intrinsic::spv_subgroup_max_size:
668 case Intrinsic::spv_subgroup_size:
669 case Intrinsic::spv_thread_id:
670 case Intrinsic::spv_thread_id_in_group:
671 case Intrinsic::spv_udot:
672 case Intrinsic::spv_undef:
673 case Intrinsic::spv_value_md:
674 case Intrinsic::spv_workgroup_size:
686 case SPIRV::OpTypeVoid:
687 case SPIRV::OpTypeBool:
688 case SPIRV::OpTypeInt:
689 case SPIRV::OpTypeFloat:
690 case SPIRV::OpTypeVector:
691 case SPIRV::OpTypeMatrix:
692 case SPIRV::OpTypeImage:
693 case SPIRV::OpTypeSampler:
694 case SPIRV::OpTypeSampledImage:
695 case SPIRV::OpTypeArray:
696 case SPIRV::OpTypeRuntimeArray:
697 case SPIRV::OpTypeStruct:
698 case SPIRV::OpTypeOpaque:
699 case SPIRV::OpTypePointer:
700 case SPIRV::OpTypeFunction:
701 case SPIRV::OpTypeEvent:
702 case SPIRV::OpTypeDeviceEvent:
703 case SPIRV::OpTypeReserveId:
704 case SPIRV::OpTypeQueue:
705 case SPIRV::OpTypePipe:
706 case SPIRV::OpTypeForwardPointer:
707 case SPIRV::OpTypePipeStorage:
708 case SPIRV::OpTypeNamedBarrier:
709 case SPIRV::OpTypeAccelerationStructureNV:
710 case SPIRV::OpTypeCooperativeMatrixNV:
711 case SPIRV::OpTypeCooperativeMatrixKHR:
721 if (
MI.getNumDefs() == 0)
724 for (
const auto &MO :
MI.all_defs()) {
726 if (
Reg.isPhysical()) {
731 if (
UseMI.getOpcode() != SPIRV::OpName) {
738 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
739 MI.isLifetimeMarker()) {
742 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
753 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
754 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
757 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
762 if (
MI.mayStore() ||
MI.isCall() ||
763 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
764 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
765 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
776 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
783void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
785 for (
const auto &MO :
MI.all_defs()) {
789 SmallVector<MachineInstr *, 4> UselessOpNames;
792 "There is still a use of the dead function.");
795 for (MachineInstr *OpNameMI : UselessOpNames) {
797 OpNameMI->eraseFromParent();
802void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
805 removeOpNamesForDeadMI(
MI);
806 MI.eraseFromParent();
809bool SPIRVInstructionSelector::select(MachineInstr &
I) {
810 resetVRegsType(*
I.getParent()->getParent());
812 assert(
I.getParent() &&
"Instruction should be in a basic block!");
813 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
818 removeDeadInstruction(
I);
825 if (Opcode == SPIRV::ASSIGN_TYPE) {
826 Register DstReg =
I.getOperand(0).getReg();
827 Register SrcReg =
I.getOperand(1).getReg();
830 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
831 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
832 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
833 Register SelectDstReg =
Def->getOperand(0).getReg();
834 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
836 assert(SuccessToSelectSelect);
838 Def->eraseFromParent();
845 bool Res = selectImpl(
I, *CoverageInfo);
847 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
848 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
852 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
864 }
else if (
I.getNumDefs() == 1) {
876 removeDeadInstruction(
I);
881 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
882 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
888 bool HasDefs =
I.getNumDefs() > 0;
891 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
892 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
893 if (spvSelect(ResVReg, ResType,
I)) {
895 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
906 case TargetOpcode::G_CONSTANT:
907 case TargetOpcode::G_FCONSTANT:
914 MachineInstr &
I)
const {
917 if (DstRC != SrcRC && SrcRC)
919 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
926bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
927 SPIRVTypeInst ResType,
928 MachineInstr &
I)
const {
929 const unsigned Opcode =
I.getOpcode();
931 return selectImpl(
I, *CoverageInfo);
933 case TargetOpcode::G_CONSTANT:
934 case TargetOpcode::G_FCONSTANT:
935 return selectConst(ResVReg, ResType,
I);
936 case TargetOpcode::G_GLOBAL_VALUE:
937 return selectGlobalValue(ResVReg,
I);
938 case TargetOpcode::G_IMPLICIT_DEF:
939 return selectOpUndef(ResVReg, ResType,
I);
940 case TargetOpcode::G_FREEZE:
941 return selectFreeze(ResVReg, ResType,
I);
943 case TargetOpcode::G_INTRINSIC:
944 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
945 case TargetOpcode::G_INTRINSIC_CONVERGENT:
946 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
947 return selectIntrinsic(ResVReg, ResType,
I);
948 case TargetOpcode::G_BITREVERSE:
949 return selectBitreverse(ResVReg, ResType,
I);
951 case TargetOpcode::G_BUILD_VECTOR:
952 return selectBuildVector(ResVReg, ResType,
I);
953 case TargetOpcode::G_SPLAT_VECTOR:
954 return selectSplatVector(ResVReg, ResType,
I);
956 case TargetOpcode::G_SHUFFLE_VECTOR: {
957 MachineBasicBlock &BB = *
I.getParent();
958 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
961 .
addUse(
I.getOperand(1).getReg())
962 .
addUse(
I.getOperand(2).getReg());
963 for (
auto V :
I.getOperand(3).getShuffleMask())
968 case TargetOpcode::G_MEMMOVE:
969 case TargetOpcode::G_MEMCPY:
970 case TargetOpcode::G_MEMSET:
971 return selectMemOperation(ResVReg,
I);
973 case TargetOpcode::G_ICMP:
974 return selectICmp(ResVReg, ResType,
I);
975 case TargetOpcode::G_FCMP:
976 return selectFCmp(ResVReg, ResType,
I);
978 case TargetOpcode::G_FRAME_INDEX:
979 return selectFrameIndex(ResVReg, ResType,
I);
981 case TargetOpcode::G_LOAD:
982 return selectLoad(ResVReg, ResType,
I);
983 case TargetOpcode::G_STORE:
984 return selectStore(
I);
986 case TargetOpcode::G_BR:
987 return selectBranch(
I);
988 case TargetOpcode::G_BRCOND:
989 return selectBranchCond(
I);
991 case TargetOpcode::G_PHI:
992 return selectPhi(ResVReg,
I);
994 case TargetOpcode::G_FPTOSI:
995 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
996 case TargetOpcode::G_FPTOUI:
997 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
999 case TargetOpcode::G_FPTOSI_SAT:
1000 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1001 case TargetOpcode::G_FPTOUI_SAT:
1002 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1004 case TargetOpcode::G_SITOFP:
1005 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1006 case TargetOpcode::G_UITOFP:
1007 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1009 case TargetOpcode::G_CTPOP:
1010 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitCount);
1011 case TargetOpcode::G_SMIN:
1012 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1013 case TargetOpcode::G_UMIN:
1014 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1016 case TargetOpcode::G_SMAX:
1017 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1018 case TargetOpcode::G_UMAX:
1019 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1021 case TargetOpcode::G_SCMP:
1022 return selectSUCmp(ResVReg, ResType,
I,
true);
1023 case TargetOpcode::G_UCMP:
1024 return selectSUCmp(ResVReg, ResType,
I,
false);
1025 case TargetOpcode::G_LROUND:
1026 case TargetOpcode::G_LLROUND: {
1029 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1031 regForLround, *(
I.getParent()->getParent()));
1033 CL::round, GL::Round,
false);
1035 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1042 case TargetOpcode::G_STRICT_FMA:
1043 case TargetOpcode::G_FMA: {
1046 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1049 .
addUse(
I.getOperand(1).getReg())
1050 .
addUse(
I.getOperand(2).getReg())
1051 .
addUse(
I.getOperand(3).getReg())
1056 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1059 case TargetOpcode::G_STRICT_FLDEXP:
1060 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1062 case TargetOpcode::G_FPOW:
1063 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1064 case TargetOpcode::G_FPOWI:
1065 return selectFpowi(ResVReg, ResType,
I);
1067 case TargetOpcode::G_FEXP:
1068 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1069 case TargetOpcode::G_FEXP2:
1070 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1071 case TargetOpcode::G_FEXP10:
1072 return selectExp10(ResVReg, ResType,
I);
1074 case TargetOpcode::G_FMODF:
1075 return selectModf(ResVReg, ResType,
I);
1076 case TargetOpcode::G_FSINCOS:
1077 return selectSincos(ResVReg, ResType,
I);
1079 case TargetOpcode::G_FLOG:
1080 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1081 case TargetOpcode::G_FLOG2:
1082 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1083 case TargetOpcode::G_FLOG10:
1084 return selectLog10(ResVReg, ResType,
I);
1086 case TargetOpcode::G_FABS:
1087 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1088 case TargetOpcode::G_ABS:
1089 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1091 case TargetOpcode::G_FMINNUM:
1092 case TargetOpcode::G_FMINIMUM:
1093 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1094 case TargetOpcode::G_FMAXNUM:
1095 case TargetOpcode::G_FMAXIMUM:
1096 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1098 case TargetOpcode::G_FCOPYSIGN:
1099 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1101 case TargetOpcode::G_FCEIL:
1102 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1103 case TargetOpcode::G_FFLOOR:
1104 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1106 case TargetOpcode::G_FCOS:
1107 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1108 case TargetOpcode::G_FSIN:
1109 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1110 case TargetOpcode::G_FTAN:
1111 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1112 case TargetOpcode::G_FACOS:
1113 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1114 case TargetOpcode::G_FASIN:
1115 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1116 case TargetOpcode::G_FATAN:
1117 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1118 case TargetOpcode::G_FATAN2:
1119 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1120 case TargetOpcode::G_FCOSH:
1121 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1122 case TargetOpcode::G_FSINH:
1123 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1124 case TargetOpcode::G_FTANH:
1125 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1127 case TargetOpcode::G_STRICT_FSQRT:
1128 case TargetOpcode::G_FSQRT:
1129 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1131 case TargetOpcode::G_CTTZ:
1132 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1133 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1134 case TargetOpcode::G_CTLZ:
1135 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1136 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1138 case TargetOpcode::G_INTRINSIC_ROUND:
1139 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1140 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1141 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1142 case TargetOpcode::G_INTRINSIC_TRUNC:
1143 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1144 case TargetOpcode::G_FRINT:
1145 case TargetOpcode::G_FNEARBYINT:
1146 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1148 case TargetOpcode::G_SMULH:
1149 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1150 case TargetOpcode::G_UMULH:
1151 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1153 case TargetOpcode::G_SADDSAT:
1154 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1155 case TargetOpcode::G_UADDSAT:
1156 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1157 case TargetOpcode::G_SSUBSAT:
1158 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1159 case TargetOpcode::G_USUBSAT:
1160 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1162 case TargetOpcode::G_FFREXP:
1163 return selectFrexp(ResVReg, ResType,
I);
1165 case TargetOpcode::G_UADDO:
1166 return selectOverflowArith(ResVReg, ResType,
I,
1167 ResType->
getOpcode() == SPIRV::OpTypeVector
1168 ? SPIRV::OpIAddCarryV
1169 : SPIRV::OpIAddCarryS);
1170 case TargetOpcode::G_USUBO:
1171 return selectOverflowArith(ResVReg, ResType,
I,
1172 ResType->
getOpcode() == SPIRV::OpTypeVector
1173 ? SPIRV::OpISubBorrowV
1174 : SPIRV::OpISubBorrowS);
1175 case TargetOpcode::G_UMULO:
1176 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1177 case TargetOpcode::G_SMULO:
1178 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1180 case TargetOpcode::G_SEXT:
1181 return selectExt(ResVReg, ResType,
I,
true);
1182 case TargetOpcode::G_ANYEXT:
1183 case TargetOpcode::G_ZEXT:
1184 return selectExt(ResVReg, ResType,
I,
false);
1185 case TargetOpcode::G_TRUNC:
1186 return selectTrunc(ResVReg, ResType,
I);
1187 case TargetOpcode::G_FPTRUNC:
1188 case TargetOpcode::G_FPEXT:
1189 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1191 case TargetOpcode::G_PTRTOINT:
1192 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1193 case TargetOpcode::G_INTTOPTR:
1194 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1195 case TargetOpcode::G_BITCAST:
1196 return selectBitcast(ResVReg, ResType,
I);
1197 case TargetOpcode::G_ADDRSPACE_CAST:
1198 return selectAddrSpaceCast(ResVReg, ResType,
I);
1199 case TargetOpcode::G_PTR_ADD: {
1201 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1205 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1206 (*II).getOpcode() == TargetOpcode::COPY ||
1207 (*II).getOpcode() == SPIRV::OpVariable) &&
1208 getImm(
I.getOperand(2), MRI));
1210 bool IsGVInit =
false;
1214 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1215 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1216 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1217 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1227 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1240 "incompatible result and operand types in a bitcast");
1242 MachineInstrBuilder MIB =
1243 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1250 : SPIRV::OpInBoundsPtrAccessChain))
1254 .
addUse(
I.getOperand(2).getReg())
1257 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1261 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1263 .
addUse(
I.getOperand(2).getReg())
1272 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1275 .
addImm(
static_cast<uint32_t
>(
1276 SPIRV::Opcode::InBoundsPtrAccessChain))
1279 .
addUse(
I.getOperand(2).getReg());
1284 case TargetOpcode::G_ATOMICRMW_OR:
1285 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1286 case TargetOpcode::G_ATOMICRMW_ADD:
1287 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1288 case TargetOpcode::G_ATOMICRMW_AND:
1289 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1290 case TargetOpcode::G_ATOMICRMW_MAX:
1291 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1292 case TargetOpcode::G_ATOMICRMW_MIN:
1293 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1294 case TargetOpcode::G_ATOMICRMW_SUB:
1295 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1296 case TargetOpcode::G_ATOMICRMW_XOR:
1297 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1298 case TargetOpcode::G_ATOMICRMW_UMAX:
1299 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1300 case TargetOpcode::G_ATOMICRMW_UMIN:
1301 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1302 case TargetOpcode::G_ATOMICRMW_XCHG:
1303 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1304 case TargetOpcode::G_ATOMIC_CMPXCHG:
1305 return selectAtomicCmpXchg(ResVReg, ResType,
I);
1307 case TargetOpcode::G_ATOMICRMW_FADD:
1308 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1309 case TargetOpcode::G_ATOMICRMW_FSUB:
1311 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1312 ResType->
getOpcode() == SPIRV::OpTypeVector
1314 : SPIRV::OpFNegate);
1315 case TargetOpcode::G_ATOMICRMW_FMIN:
1316 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1317 case TargetOpcode::G_ATOMICRMW_FMAX:
1318 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1320 case TargetOpcode::G_FENCE:
1321 return selectFence(
I);
1323 case TargetOpcode::G_STACKSAVE:
1324 return selectStackSave(ResVReg, ResType,
I);
1325 case TargetOpcode::G_STACKRESTORE:
1326 return selectStackRestore(
I);
1328 case TargetOpcode::G_UNMERGE_VALUES:
1334 case TargetOpcode::G_TRAP:
1335 case TargetOpcode::G_UBSANTRAP:
1336 case TargetOpcode::DBG_LABEL:
1338 case TargetOpcode::G_DEBUGTRAP:
1339 return selectDebugTrap(ResVReg, ResType,
I);
1346bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1347 SPIRVTypeInst ResType,
1348 MachineInstr &
I)
const {
1349 unsigned Opcode = SPIRV::OpNop;
1356bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1357 SPIRVTypeInst ResType,
1359 GL::GLSLExtInst GLInst,
1360 bool setMIFlags,
bool useMISrc,
1363 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1364 std::string DiagMsg;
1365 raw_string_ostream OS(DiagMsg);
1366 I.print(OS,
true,
false,
false,
false);
1367 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1370 return selectExtInst(ResVReg, ResType,
I,
1371 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1372 setMIFlags, useMISrc, SrcRegs);
1375bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1376 SPIRVTypeInst ResType,
1378 CL::OpenCLExtInst CLInst,
1379 bool setMIFlags,
bool useMISrc,
1381 return selectExtInst(ResVReg, ResType,
I,
1382 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1383 setMIFlags, useMISrc, SrcRegs);
1386bool SPIRVInstructionSelector::selectExtInst(
1387 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1388 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1390 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1391 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1392 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1396bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1397 SPIRVTypeInst ResType,
1400 bool setMIFlags,
bool useMISrc,
1403 for (
const auto &[InstructionSet, Opcode] : Insts) {
1407 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1410 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1415 const unsigned NumOps =
I.getNumOperands();
1418 I.getOperand(Index).getType() ==
1419 MachineOperand::MachineOperandType::MO_IntrinsicID)
1422 MIB.
add(
I.getOperand(Index));
1434bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1435 SPIRVTypeInst ResType,
1436 MachineInstr &
I)
const {
1437 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1438 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1439 for (
const auto &Ex : ExtInsts) {
1440 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1441 uint32_t Opcode = Ex.second;
1445 MachineIRBuilder MIRBuilder(
I);
1448 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1453 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1456 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1459 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1462 .
addImm(
static_cast<uint32_t
>(Ex.first))
1464 .
add(
I.getOperand(2))
1468 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1469 .
addDef(
I.getOperand(1).getReg())
1478bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1479 SPIRVTypeInst ResType,
1480 MachineInstr &
I)
const {
1481 Register CosResVReg =
I.getOperand(1).getReg();
1482 unsigned SrcIdx =
I.getNumExplicitDefs();
1487 MachineIRBuilder MIRBuilder(
I);
1489 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1494 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1497 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1499 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1502 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1504 .
add(
I.getOperand(SrcIdx))
1507 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1515 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1518 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1520 .
add(
I.getOperand(SrcIdx))
1522 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1525 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1527 .
add(
I.getOperand(SrcIdx))
1534bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1535 SPIRVTypeInst ResType,
1537 std::vector<Register> Srcs,
1538 unsigned Opcode)
const {
1539 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1549bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1550 SPIRVTypeInst ResType,
1552 unsigned Opcode)
const {
1554 Register SrcReg =
I.getOperand(1).getReg();
1559 unsigned DefOpCode = DefIt->getOpcode();
1560 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1563 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1564 DefOpCode = VRD->getOpcode();
1566 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1567 DefOpCode == TargetOpcode::G_CONSTANT ||
1568 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1574 uint32_t SpecOpcode = 0;
1576 case SPIRV::OpConvertPtrToU:
1577 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1579 case SPIRV::OpConvertUToPtr:
1580 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1585 TII.get(SPIRV::OpSpecConstantOp))
1595 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1599bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1600 SPIRVTypeInst ResType,
1601 MachineInstr &
I)
const {
1602 Register OpReg =
I.getOperand(1).getReg();
1603 SPIRVTypeInst OpType =
1607 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1617 if (
MemOp->isVolatile())
1618 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1619 if (
MemOp->isNonTemporal())
1620 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1622 if (!ST->isShader() &&
MemOp->getAlign().value())
1623 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1627 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1628 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1632 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1634 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1638 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1642 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1644 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1656 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1658 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1660 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1664bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1665 SPIRVTypeInst ResType,
1666 MachineInstr &
I)
const {
1668 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1673 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1674 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1676 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1680 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1684 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1685 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1686 I.getDebugLoc(),
I);
1690 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1694 if (!
I.getNumMemOperands()) {
1695 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1697 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1700 MachineIRBuilder MIRBuilder(
I);
1707bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1709 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1710 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1715 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1716 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1721 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1725 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1726 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1727 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1728 TII.get(SPIRV::OpImageWrite))
1734 if (sampledTypeIsSignedInteger(LLVMHandleType))
1737 BMI.constrainAllUses(
TII,
TRI, RBI);
1743 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpStore))
1746 if (!
I.getNumMemOperands()) {
1747 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1749 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1752 MachineIRBuilder MIRBuilder(
I);
1759bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
1760 SPIRVTypeInst ResType,
1761 MachineInstr &
I)
const {
1762 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
1770 const Register PtrsReg =
I.getOperand(2).getReg();
1771 const uint32_t Alignment =
I.getOperand(3).getImm();
1772 const Register MaskReg =
I.getOperand(4).getReg();
1773 const Register PassthruReg =
I.getOperand(5).getReg();
1774 const Register AlignmentReg = buildI32Constant(Alignment,
I);
1778 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
1789bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
1790 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
1797 const Register ValuesReg =
I.getOperand(1).getReg();
1798 const Register PtrsReg =
I.getOperand(2).getReg();
1799 const uint32_t Alignment =
I.getOperand(3).getImm();
1800 const Register MaskReg =
I.getOperand(4).getReg();
1801 const Register AlignmentReg = buildI32Constant(Alignment,
I);
1805 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
1814bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
1815 const Twine &Msg)
const {
1816 const Function &
F =
I.getMF()->getFunction();
1817 F.getContext().diagnose(
1818 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
1822bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
1823 SPIRVTypeInst ResType,
1824 MachineInstr &
I)
const {
1825 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1827 "llvm.stacksave intrinsic: this instruction requires the following "
1828 "SPIR-V extension: SPV_INTEL_variable_length_array",
1831 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
1838bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
1839 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1841 "llvm.stackrestore intrinsic: this instruction requires the following "
1842 "SPIR-V extension: SPV_INTEL_variable_length_array",
1844 if (!
I.getOperand(0).isReg())
1847 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
1848 .
addUse(
I.getOperand(0).getReg())
1854SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
1855 MachineIRBuilder MIRBuilder(
I);
1856 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1863 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
1867 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1868 Type *ArrTy = ArrayType::get(ValTy, Num);
1870 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
1873 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
1880 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
1883 .
addImm(SPIRV::StorageClass::UniformConstant)
1894bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
1897 Register DstReg =
I.getOperand(0).getReg();
1902 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
1907 "Unable to determine pointee type size for OpCopyMemory");
1908 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
1909 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
1911 "OpCopyMemory requires the size to match the pointee type size");
1912 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
1915 if (
I.getNumMemOperands()) {
1916 MachineIRBuilder MIRBuilder(
I);
1923bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
1926 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
1927 .
addUse(
I.getOperand(0).getReg())
1929 .
addUse(
I.getOperand(2).getReg());
1930 if (
I.getNumMemOperands()) {
1931 MachineIRBuilder MIRBuilder(
I);
1938bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
1939 MachineInstr &
I)
const {
1940 Register SrcReg =
I.getOperand(1).getReg();
1941 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
1942 Register VarReg = getOrCreateMemSetGlobal(
I);
1945 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1947 ValTy,
I, SPIRV::StorageClass::UniformConstant);
1949 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
1953 if (!selectCopyMemory(
I, SrcReg))
1956 if (!selectCopyMemorySized(
I, SrcReg))
1959 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
1960 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
1965bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
1966 SPIRVTypeInst ResType,
1969 unsigned NegateOpcode)
const {
1971 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1974 Register ScopeReg = buildI32Constant(Scope,
I);
1976 Register Ptr =
I.getOperand(1).getReg();
1982 Register MemSemReg = buildI32Constant(MemSem ,
I);
1984 Register ValueReg =
I.getOperand(2).getReg();
1985 if (NegateOpcode != 0) {
1988 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
1993 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2004bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2005 unsigned ArgI =
I.getNumOperands() - 1;
2007 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2008 SPIRVTypeInst SrcType =
2010 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2012 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2014 SPIRVTypeInst ScalarType =
2017 unsigned CurrentIndex = 0;
2018 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2019 Register ResVReg =
I.getOperand(i).getReg();
2022 LLT ResLLT = MRI->
getType(ResVReg);
2028 ResType = ScalarType;
2034 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2037 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2043 for (
unsigned j = 0;
j < NumElements; ++
j) {
2044 MIB.
addImm(CurrentIndex + j);
2046 CurrentIndex += NumElements;
2050 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2062bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2065 Register MemSemReg = buildI32Constant(MemSem,
I);
2067 uint32_t
Scope =
static_cast<uint32_t
>(
2069 Register ScopeReg = buildI32Constant(Scope,
I);
2071 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2078bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2079 SPIRVTypeInst ResType,
2081 unsigned Opcode)
const {
2082 Type *ResTy =
nullptr;
2086 "Not enough info to select the arithmetic with overflow instruction");
2089 "with overflow instruction");
2095 MachineIRBuilder MIRBuilder(
I);
2097 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2098 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2104 Register ZeroReg = buildZerosVal(ResType,
I);
2109 if (ResName.
size() > 0)
2114 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2117 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2118 MIB.
addUse(
I.getOperand(i).getReg());
2123 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2124 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2126 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2127 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2134 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2135 .
addDef(
I.getOperand(1).getReg())
2143bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2144 SPIRVTypeInst ResType,
2145 MachineInstr &
I)
const {
2149 Register Ptr =
I.getOperand(2).getReg();
2152 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2155 ScopeReg = buildI32Constant(Scope,
I);
2157 unsigned ScSem =
static_cast<uint32_t
>(
2160 unsigned MemSemEq =
static_cast<uint32_t
>(
getMemSemantics(AO)) | ScSem;
2161 Register MemSemEqReg = buildI32Constant(MemSemEq,
I);
2163 unsigned MemSemNeq =
static_cast<uint32_t
>(
getMemSemantics(FO)) | ScSem;
2164 if (MemSemEq == MemSemNeq)
2165 MemSemNeqReg = MemSemEqReg;
2167 MemSemNeqReg = buildI32Constant(MemSemEq,
I);
2170 ScopeReg =
I.getOperand(5).getReg();
2171 MemSemEqReg =
I.getOperand(6).getReg();
2172 MemSemNeqReg =
I.getOperand(7).getReg();
2176 Register Val =
I.getOperand(4).getReg();
2180 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2199 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2206 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2218 case SPIRV::StorageClass::DeviceOnlyINTEL:
2219 case SPIRV::StorageClass::HostOnlyINTEL:
2228 bool IsGRef =
false;
2229 bool IsAllowedRefs =
2231 unsigned Opcode = It.getOpcode();
2232 if (Opcode == SPIRV::OpConstantComposite ||
2233 Opcode == SPIRV::OpSpecConstantComposite ||
2234 Opcode == SPIRV::OpVariable ||
2235 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2236 return IsGRef = true;
2237 return Opcode == SPIRV::OpName;
2239 return IsAllowedRefs && IsGRef;
2242Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2243 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2245 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2249SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2251 uint32_t Opcode)
const {
2252 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2253 TII.get(SPIRV::OpSpecConstantOp))
2261SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2262 SPIRVTypeInst SrcPtrTy)
const {
2263 SPIRVTypeInst GenericPtrTy =
2267 SPIRV::StorageClass::Generic),
2269 MachineFunction *MF =
I.getParent()->getParent();
2271 MachineInstrBuilder MIB = buildSpecConstantOp(
2273 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2283bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2284 SPIRVTypeInst ResType,
2285 MachineInstr &
I)
const {
2289 Register SrcPtr =
I.getOperand(1).getReg();
2293 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2294 ResType->
getOpcode() != SPIRV::OpTypePointer)
2295 return BuildCOPY(ResVReg, SrcPtr,
I);
2305 unsigned SpecOpcode =
2307 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2310 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2317 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2319 .constrainAllUses(
TII,
TRI, RBI);
2321 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2323 buildSpecConstantOp(
2325 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2326 .constrainAllUses(
TII,
TRI, RBI);
2333 return BuildCOPY(ResVReg, SrcPtr,
I);
2335 if ((SrcSC == SPIRV::StorageClass::Function &&
2336 DstSC == SPIRV::StorageClass::Private) ||
2337 (DstSC == SPIRV::StorageClass::Function &&
2338 SrcSC == SPIRV::StorageClass::Private))
2339 return BuildCOPY(ResVReg, SrcPtr,
I);
2343 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2346 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2349 SPIRVTypeInst GenericPtrTy =
2368 return selectUnOp(ResVReg, ResType,
I,
2369 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2371 return selectUnOp(ResVReg, ResType,
I,
2372 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2374 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2376 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2386 return SPIRV::OpFOrdEqual;
2388 return SPIRV::OpFOrdGreaterThanEqual;
2390 return SPIRV::OpFOrdGreaterThan;
2392 return SPIRV::OpFOrdLessThanEqual;
2394 return SPIRV::OpFOrdLessThan;
2396 return SPIRV::OpFOrdNotEqual;
2398 return SPIRV::OpOrdered;
2400 return SPIRV::OpFUnordEqual;
2402 return SPIRV::OpFUnordGreaterThanEqual;
2404 return SPIRV::OpFUnordGreaterThan;
2406 return SPIRV::OpFUnordLessThanEqual;
2408 return SPIRV::OpFUnordLessThan;
2410 return SPIRV::OpFUnordNotEqual;
2412 return SPIRV::OpUnordered;
2422 return SPIRV::OpIEqual;
2424 return SPIRV::OpINotEqual;
2426 return SPIRV::OpSGreaterThanEqual;
2428 return SPIRV::OpSGreaterThan;
2430 return SPIRV::OpSLessThanEqual;
2432 return SPIRV::OpSLessThan;
2434 return SPIRV::OpUGreaterThanEqual;
2436 return SPIRV::OpUGreaterThan;
2438 return SPIRV::OpULessThanEqual;
2440 return SPIRV::OpULessThan;
2449 return SPIRV::OpPtrEqual;
2451 return SPIRV::OpPtrNotEqual;
2462 return SPIRV::OpLogicalEqual;
2464 return SPIRV::OpLogicalNotEqual;
2498bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2499 SPIRVTypeInst ResType,
2501 unsigned OpAnyOrAll)
const {
2502 assert(
I.getNumOperands() == 3);
2503 assert(
I.getOperand(2).isReg());
2505 Register InputRegister =
I.getOperand(2).getReg();
2508 assert(InputType &&
"VReg has no type assigned");
2511 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2512 if (IsBoolTy && !IsVectorTy) {
2513 assert(ResVReg ==
I.getOperand(0).getReg());
2514 return BuildCOPY(ResVReg, InputRegister,
I);
2518 unsigned SpirvNotEqualId =
2519 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2521 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2526 IsBoolTy ? InputRegister
2534 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2536 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2553bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2554 SPIRVTypeInst ResType,
2555 MachineInstr &
I)
const {
2556 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2559bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2560 SPIRVTypeInst ResType,
2561 MachineInstr &
I)
const {
2562 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2566bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2567 SPIRVTypeInst ResType,
2568 MachineInstr &
I)
const {
2569 assert(
I.getNumOperands() == 4);
2570 assert(
I.getOperand(2).isReg());
2571 assert(
I.getOperand(3).isReg());
2573 [[maybe_unused]] SPIRVTypeInst VecType =
2578 "dot product requires a vector of at least 2 components");
2580 [[maybe_unused]] SPIRVTypeInst EltType =
2589 .
addUse(
I.getOperand(2).getReg())
2590 .
addUse(
I.getOperand(3).getReg())
2595bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2596 SPIRVTypeInst ResType,
2599 assert(
I.getNumOperands() == 4);
2600 assert(
I.getOperand(2).isReg());
2601 assert(
I.getOperand(3).isReg());
2604 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2608 .
addUse(
I.getOperand(2).getReg())
2609 .
addUse(
I.getOperand(3).getReg())
2616bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2617 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2618 assert(
I.getNumOperands() == 4);
2619 assert(
I.getOperand(2).isReg());
2620 assert(
I.getOperand(3).isReg());
2624 Register Vec0 =
I.getOperand(2).getReg();
2625 Register Vec1 =
I.getOperand(3).getReg();
2629 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2638 "dot product requires a vector of at least 2 components");
2641 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2651 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2662 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2674bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2675 SPIRVTypeInst ResType,
2676 MachineInstr &
I)
const {
2678 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2681 .
addUse(
I.getOperand(2).getReg())
2686bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2687 SPIRVTypeInst ResType,
2688 MachineInstr &
I)
const {
2690 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2693 .
addUse(
I.getOperand(2).getReg())
2698template <
bool Signed>
2699bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2700 SPIRVTypeInst ResType,
2701 MachineInstr &
I)
const {
2702 assert(
I.getNumOperands() == 5);
2703 assert(
I.getOperand(2).isReg());
2704 assert(
I.getOperand(3).isReg());
2705 assert(
I.getOperand(4).isReg());
2708 Register Acc =
I.getOperand(2).getReg();
2712 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2714 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
2719 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
2722 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2734template <
bool Signed>
2735bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
2736 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2737 assert(
I.getNumOperands() == 5);
2738 assert(
I.getOperand(2).isReg());
2739 assert(
I.getOperand(3).isReg());
2740 assert(
I.getOperand(4).isReg());
2743 Register Acc =
I.getOperand(2).getReg();
2749 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
2753 for (
unsigned i = 0; i < 4; i++) {
2776 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
2796 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2811bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
2812 SPIRVTypeInst ResType,
2813 MachineInstr &
I)
const {
2814 assert(
I.getNumOperands() == 3);
2815 assert(
I.getOperand(2).isReg());
2817 Register VZero = buildZerosValF(ResType,
I);
2818 Register VOne = buildOnesValF(ResType,
I);
2820 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
2823 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2825 .
addUse(
I.getOperand(2).getReg())
2832bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
2833 SPIRVTypeInst ResType,
2834 MachineInstr &
I)
const {
2835 assert(
I.getNumOperands() == 3);
2836 assert(
I.getOperand(2).isReg());
2838 Register InputRegister =
I.getOperand(2).getReg();
2840 auto &
DL =
I.getDebugLoc();
2850 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
2852 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
2860 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2865 if (NeedsConversion) {
2866 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
2877bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
2878 SPIRVTypeInst ResType,
2880 unsigned Opcode)
const {
2884 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2890 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
2891 BMI.addUse(
I.getOperand(J).getReg());
2898bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
2901 bool WithGroupSync)
const {
2903 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
2905 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
2907 assert(((Scope != SPIRV::Scope::Workgroup) ||
2908 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
2909 "Workgroup Scope must set WorkGroupMemory semantic "
2910 "in Barrier instruction");
2912 assert(((Scope != SPIRV::Scope::Device) ||
2913 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
2914 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
2915 "Device Scope must set UniformMemory and ImageMemory semantic "
2916 "in Barrier instruction");
2918 Register MemSemReg = buildI32Constant(MemSem,
I);
2919 Register ScopeReg = buildI32Constant(Scope,
I);
2925 if (WithGroupSync) {
2926 MI.addUse(ScopeReg);
2929 MI.addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
2933bool SPIRVInstructionSelector::selectWaveActiveCountBits(
2934 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2939 if (!selectWaveOpInst(BallotReg, BallotType,
I,
2940 SPIRV::OpGroupNonUniformBallot))
2945 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2950 .
addImm(SPIRV::GroupOperation::Reduce)
2959 if (
Type->getOpcode() != SPIRV::OpTypeVector)
2963 return Type->getOperand(2).getImm();
2966bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
2967 SPIRVTypeInst ResType,
2968 MachineInstr &
I)
const {
2973 Register InputReg =
I.getOperand(2).getReg();
2978 bool IsVector = NumElems > 1;
2981 SPIRVTypeInst ElemInputType = InputType;
2982 SPIRVTypeInst ElemBoolType = ResType;
2995 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
2996 SPIRV::OpGroupNonUniformAllEqual);
3001 ElementResults.
reserve(NumElems);
3003 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3016 ElemInput = Extracted;
3022 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3033 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3044bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3045 SPIRVTypeInst ResType,
3046 MachineInstr &
I)
const {
3048 assert(
I.getNumOperands() == 3);
3050 auto Op =
I.getOperand(2);
3062 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3084 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3088 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3095bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3096 SPIRVTypeInst ResType,
3098 bool IsUnsigned)
const {
3099 return selectWaveReduce(
3100 ResVReg, ResType,
I, IsUnsigned,
3101 [&](
Register InputRegister,
bool IsUnsigned) {
3102 const bool IsFloatTy =
3104 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3105 : SPIRV::OpGroupNonUniformSMax;
3106 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3110bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3111 SPIRVTypeInst ResType,
3113 bool IsUnsigned)
const {
3114 return selectWaveReduce(
3115 ResVReg, ResType,
I, IsUnsigned,
3116 [&](
Register InputRegister,
bool IsUnsigned) {
3117 const bool IsFloatTy =
3119 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3120 : SPIRV::OpGroupNonUniformSMin;
3121 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3125bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3126 SPIRVTypeInst ResType,
3127 MachineInstr &
I)
const {
3128 return selectWaveReduce(ResVReg, ResType,
I,
false,
3129 [&](
Register InputRegister,
bool IsUnsigned) {
3131 InputRegister, SPIRV::OpTypeFloat);
3132 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3133 : SPIRV::OpGroupNonUniformIAdd;
3137bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3138 SPIRVTypeInst ResType,
3139 MachineInstr &
I)
const {
3140 return selectWaveReduce(ResVReg, ResType,
I,
false,
3141 [&](
Register InputRegister,
bool IsUnsigned) {
3143 InputRegister, SPIRV::OpTypeFloat);
3144 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3145 : SPIRV::OpGroupNonUniformIMul;
3149template <
typename PickOpcodeFn>
3150bool SPIRVInstructionSelector::selectWaveReduce(
3151 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3152 PickOpcodeFn &&PickOpcode)
const {
3153 assert(
I.getNumOperands() == 3);
3154 assert(
I.getOperand(2).isReg());
3156 Register InputRegister =
I.getOperand(2).getReg();
3163 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3169 .
addImm(SPIRV::GroupOperation::Reduce)
3170 .
addUse(
I.getOperand(2).getReg())
3175bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3176 SPIRVTypeInst ResType,
3178 unsigned Opcode)
const {
3179 return selectWaveReduce(
3180 ResVReg, ResType,
I,
false,
3181 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3184bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3185 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3186 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3187 [&](
Register InputRegister,
bool IsUnsigned) {
3189 InputRegister, SPIRV::OpTypeFloat);
3191 ? SPIRV::OpGroupNonUniformFAdd
3192 : SPIRV::OpGroupNonUniformIAdd;
3196bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3197 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3198 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3199 [&](
Register InputRegister,
bool IsUnsigned) {
3201 InputRegister, SPIRV::OpTypeFloat);
3203 ? SPIRV::OpGroupNonUniformFMul
3204 : SPIRV::OpGroupNonUniformIMul;
3208template <
typename PickOpcodeFn>
3209bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3210 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3211 PickOpcodeFn &&PickOpcode)
const {
3212 assert(
I.getNumOperands() == 3);
3213 assert(
I.getOperand(2).isReg());
3215 Register InputRegister =
I.getOperand(2).getReg();
3222 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3228 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3229 .
addUse(
I.getOperand(2).getReg())
3234bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3235 SPIRVTypeInst ResType,
3238 assert(
I.getNumOperands() == 3);
3239 assert(
I.getOperand(2).isReg());
3241 Register InputRegister =
I.getOperand(2).getReg();
3247 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3258bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3259 SPIRVTypeInst ResType,
3264 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3269 : SPIRV::OpUConvert;
3273 ShiftOp = SPIRV::OpShiftRightLogicalV;
3278 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3279 TII.get(SPIRV::OpConstantComposite))
3282 for (
unsigned It = 0; It <
N; ++It)
3286 ShiftConst = CompositeReg;
3291 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3296 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3301 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3306 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3309bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
3310 SPIRVTypeInst ResType,
3314 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3322bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3323 SPIRVTypeInst ResType,
3324 MachineInstr &
I)
const {
3325 Register OpReg =
I.getOperand(1).getReg();
3332 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3334 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3339 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3350 unsigned AndOp = SPIRV::OpBitwiseAndS;
3351 unsigned OrOp = SPIRV::OpBitwiseOrS;
3352 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
3353 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
3355 AndOp = SPIRV::OpBitwiseAndV;
3356 OrOp = SPIRV::OpBitwiseOrV;
3357 ShlOp = SPIRV::OpShiftLeftLogicalV;
3358 ShrOp = SPIRV::OpShiftRightLogicalV;
3364 const unsigned Shift) ->
Register {
3372 Register MaskReg = CreateConst(Mask);
3373 Register ShiftReg = CreateConst(Shift);
3380 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
3381 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
3382 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
3383 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
3384 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
3392 uint64_t
Mask = ~0ull;
3393 while ((Shift >>= 1) > 0) {
3400 return BuildCOPY(ResVReg, Result,
I);
3403bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3404 SPIRVTypeInst ResType,
3405 MachineInstr &
I)
const {
3411 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3413 Register OpReg =
I.getOperand(1).getReg();
3414 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
3415 if (
Def->getOpcode() == TargetOpcode::COPY)
3418 switch (
Def->getOpcode()) {
3419 case SPIRV::ASSIGN_TYPE:
3420 if (MachineInstr *AssignToDef =
3422 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3423 Reg =
Def->getOperand(2).getReg();
3426 case SPIRV::OpUndef:
3427 Reg =
Def->getOperand(1).getReg();
3430 unsigned DestOpCode;
3432 DestOpCode = SPIRV::OpConstantNull;
3434 DestOpCode = TargetOpcode::COPY;
3437 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3438 .
addDef(
I.getOperand(0).getReg())
3446bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3447 SPIRVTypeInst ResType,
3448 MachineInstr &
I)
const {
3450 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3452 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3456 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3461 for (
unsigned i =
I.getNumExplicitDefs();
3462 i <
I.getNumExplicitOperands() && IsConst; ++i)
3466 if (!IsConst &&
N < 2)
3468 "There must be at least two constituent operands in a vector");
3471 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3472 TII.get(IsConst ? SPIRV::OpConstantComposite
3473 : SPIRV::OpCompositeConstruct))
3476 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3477 MIB.
addUse(
I.getOperand(i).getReg());
3482bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3483 SPIRVTypeInst ResType,
3484 MachineInstr &
I)
const {
3486 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3488 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3494 if (!
I.getOperand(
OpIdx).isReg())
3501 if (!IsConst &&
N < 2)
3503 "There must be at least two constituent operands in a vector");
3506 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3507 TII.get(IsConst ? SPIRV::OpConstantComposite
3508 : SPIRV::OpCompositeConstruct))
3511 for (
unsigned i = 0; i <
N; ++i)
3517bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3518 SPIRVTypeInst ResType,
3519 MachineInstr &
I)
const {
3524 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3526 Opcode = SPIRV::OpDemoteToHelperInvocation;
3528 Opcode = SPIRV::OpKill;
3530 if (MachineInstr *NextI =
I.getNextNode()) {
3532 NextI->eraseFromParent();
3542bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3543 SPIRVTypeInst ResType,
unsigned CmpOpc,
3544 MachineInstr &
I)
const {
3545 Register Cmp0 =
I.getOperand(2).getReg();
3546 Register Cmp1 =
I.getOperand(3).getReg();
3549 "CMP operands should have the same type");
3550 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3560bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3561 SPIRVTypeInst ResType,
3562 MachineInstr &
I)
const {
3563 auto Pred =
I.getOperand(1).getPredicate();
3566 Register CmpOperand =
I.getOperand(2).getReg();
3573 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3577SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3578 SPIRVTypeInst ResType)
const {
3580 SPIRVTypeInst SpvI32Ty =
3583 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3590 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3593 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3596 .
addImm(APInt(32, Val).getZExtValue());
3598 GR.
add(ConstInt,
MI);
3603bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
3604 SPIRVTypeInst ResType,
3605 MachineInstr &
I)
const {
3607 return selectCmp(ResVReg, ResType, CmpOp,
I);
3610bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
3611 SPIRVTypeInst ResType,
3612 MachineInstr &
I)
const {
3614 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
3621 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
3622 ResType->
getOpcode() != SPIRV::OpTypeFloat)
3625 MachineIRBuilder MIRBuilder(
I);
3627 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
3633 "only float operands supported by GLSL extended math");
3636 MIRBuilder, SpirvScalarType);
3638 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
3639 ? SPIRV::OpVectorTimesScalar
3642 if (!selectOpWithSrcs(ArgReg, ResType,
I,
3643 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
3645 if (!selectExtInst(ResVReg, ResType,
I,
3646 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
3656Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
3657 MachineInstr &
I)
const {
3660 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3665bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
3671 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
3679 if (
Def->getOpcode() == SPIRV::OpConstantNull)
3682 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
3683 Def->getOpcode() == SPIRV::OpConstantI)
3696 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
3697 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
3699 Intrinsic::spv_const_composite)) {
3700 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
3701 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
3702 if (!IsZero(
Def->getOperand(i).getReg()))
3711Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
3712 MachineInstr &
I)
const {
3716 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3721Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
3722 MachineInstr &
I)
const {
3726 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3732 SPIRVTypeInst ResType,
3733 MachineInstr &
I)
const {
3737 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3742bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
3743 SPIRVTypeInst ResType,
3744 MachineInstr &
I)
const {
3745 Register SelectFirstArg =
I.getOperand(2).getReg();
3746 Register SelectSecondArg =
I.getOperand(3).getReg();
3755 SPIRV::OpTypeVector;
3762 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
3763 }
else if (IsPtrTy) {
3764 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
3766 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
3770 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
3771 }
else if (IsPtrTy) {
3772 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
3774 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3777 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3780 .
addUse(
I.getOperand(1).getReg())
3789bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
3790 SPIRVTypeInst ResType,
3792 MachineInstr &InsertAt,
3793 bool IsSigned)
const {
3795 Register ZeroReg = buildZerosVal(ResType, InsertAt);
3796 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
3797 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
3799 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3811bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
3812 SPIRVTypeInst ResType,
3813 MachineInstr &
I,
bool IsSigned,
3814 unsigned Opcode)
const {
3815 Register SrcReg =
I.getOperand(1).getReg();
3821 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
3826 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
3828 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
3831bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
3832 SPIRVTypeInst ResType, MachineInstr &
I,
3833 bool IsSigned)
const {
3834 Register SrcReg =
I.getOperand(1).getReg();
3836 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
3840 if (ResType == SrcType)
3841 return BuildCOPY(ResVReg, SrcReg,
I);
3843 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3844 return selectUnOp(ResVReg, ResType,
I, Opcode);
3847bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
3848 SPIRVTypeInst ResType,
3850 bool IsSigned)
const {
3851 MachineIRBuilder MIRBuilder(
I);
3852 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
3867 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
3870 .
addUse(
I.getOperand(1).getReg())
3871 .
addUse(
I.getOperand(2).getReg())
3877 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
3880 .
addUse(
I.getOperand(1).getReg())
3881 .
addUse(
I.getOperand(2).getReg())
3889 unsigned SelectOpcode =
3890 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
3895 .
addUse(buildOnesVal(
true, ResType,
I))
3896 .
addUse(buildZerosVal(ResType,
I))
3903 .
addUse(buildOnesVal(
false, ResType,
I))
3908bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
3911 SPIRVTypeInst IntTy,
3912 SPIRVTypeInst BoolTy)
const {
3915 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
3916 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
3918 Register One = buildOnesVal(
false, IntTy,
I);
3926 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
3935bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
3936 SPIRVTypeInst ResType,
3937 MachineInstr &
I)
const {
3938 Register IntReg =
I.getOperand(1).getReg();
3941 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
3942 if (ArgType == ResType)
3943 return BuildCOPY(ResVReg, IntReg,
I);
3945 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3946 return selectUnOp(ResVReg, ResType,
I, Opcode);
3949bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
3950 SPIRVTypeInst ResType,
3951 MachineInstr &
I)
const {
3952 unsigned Opcode =
I.getOpcode();
3953 unsigned TpOpcode = ResType->
getOpcode();
3955 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
3956 assert(Opcode == TargetOpcode::G_CONSTANT &&
3957 I.getOperand(1).getCImm()->isZero());
3958 MachineBasicBlock &DepMBB =
I.getMF()->front();
3961 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
3968 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
3971bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
3972 SPIRVTypeInst ResType,
3973 MachineInstr &
I)
const {
3974 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3981bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
3982 SPIRVTypeInst ResType,
3983 MachineInstr &
I)
const {
3985 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
3989 .
addUse(
I.getOperand(3).getReg())
3991 .
addUse(
I.getOperand(2).getReg());
3992 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
3998bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
3999 SPIRVTypeInst ResType,
4000 MachineInstr &
I)
const {
4001 Type *MaybeResTy =
nullptr;
4006 "Expected aggregate type for extractv instruction");
4008 SPIRV::AccessQualifier::ReadWrite,
false);
4012 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4015 .
addUse(
I.getOperand(2).getReg());
4016 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4022bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4023 SPIRVTypeInst ResType,
4024 MachineInstr &
I)
const {
4025 if (
getImm(
I.getOperand(4), MRI))
4026 return selectInsertVal(ResVReg, ResType,
I);
4028 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4031 .
addUse(
I.getOperand(2).getReg())
4032 .
addUse(
I.getOperand(3).getReg())
4033 .
addUse(
I.getOperand(4).getReg())
4038bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4039 SPIRVTypeInst ResType,
4040 MachineInstr &
I)
const {
4041 if (
getImm(
I.getOperand(3), MRI))
4042 return selectExtractVal(ResVReg, ResType,
I);
4044 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4047 .
addUse(
I.getOperand(2).getReg())
4048 .
addUse(
I.getOperand(3).getReg())
4053bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4054 SPIRVTypeInst ResType,
4055 MachineInstr &
I)
const {
4056 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4062 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4063 : SPIRV::OpAccessChain)
4064 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4065 :
SPIRV::OpPtrAccessChain);
4067 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4071 .
addUse(
I.getOperand(3).getReg());
4073 (Opcode == SPIRV::OpPtrAccessChain ||
4074 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4075 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4076 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4079 const unsigned StartingIndex =
4080 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4083 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4084 Res.addUse(
I.getOperand(i).getReg());
4085 Res.constrainAllUses(
TII,
TRI, RBI);
4090bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4092 unsigned Lim =
I.getNumExplicitOperands();
4093 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4094 Register OpReg =
I.getOperand(i).getReg();
4095 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4097 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4098 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4099 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4106 MachineFunction *MF =
I.getMF();
4118 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4119 TII.get(SPIRV::OpSpecConstantOp))
4122 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4124 GR.
add(OpDefine, MIB);
4130bool SPIRVInstructionSelector::selectDerivativeInst(
4131 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4132 const unsigned DPdOpCode)
const {
4135 errorIfInstrOutsideShader(
I);
4140 Register SrcReg =
I.getOperand(2).getReg();
4145 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4148 .
addUse(
I.getOperand(2).getReg());
4150 MachineIRBuilder MIRBuilder(
I);
4153 if (componentCount != 1)
4157 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4161 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4166 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4171 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4179bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4180 SPIRVTypeInst ResType,
4181 MachineInstr &
I)
const {
4185 case Intrinsic::spv_load:
4186 return selectLoad(ResVReg, ResType,
I);
4187 case Intrinsic::spv_store:
4188 return selectStore(
I);
4189 case Intrinsic::spv_extractv:
4190 return selectExtractVal(ResVReg, ResType,
I);
4191 case Intrinsic::spv_insertv:
4192 return selectInsertVal(ResVReg, ResType,
I);
4193 case Intrinsic::spv_extractelt:
4194 return selectExtractElt(ResVReg, ResType,
I);
4195 case Intrinsic::spv_insertelt:
4196 return selectInsertElt(ResVReg, ResType,
I);
4197 case Intrinsic::spv_gep:
4198 return selectGEP(ResVReg, ResType,
I);
4199 case Intrinsic::spv_bitcast: {
4200 Register OpReg =
I.getOperand(2).getReg();
4201 SPIRVTypeInst OpType =
4205 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4207 case Intrinsic::spv_unref_global:
4208 case Intrinsic::spv_init_global: {
4209 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4214 Register GVarVReg =
MI->getOperand(0).getReg();
4215 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4220 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4222 MI->eraseFromParent();
4226 case Intrinsic::spv_undef: {
4227 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4233 case Intrinsic::spv_named_boolean_spec_constant: {
4234 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4235 : SPIRV::OpSpecConstantFalse;
4237 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4238 .
addDef(
I.getOperand(0).getReg())
4241 unsigned SpecId =
I.getOperand(2).getImm();
4243 SPIRV::Decoration::SpecId, {SpecId});
4247 case Intrinsic::spv_const_composite: {
4249 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4255 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4257 std::function<bool(
Register)> HasSpecConstOperand =
4267 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
4268 J < Def->getNumExplicitOperands(); ++J) {
4269 if (
Def->getOperand(J).isReg() &&
4270 HasSpecConstOperand(
Def->getOperand(J).getReg()))
4276 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
4277 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
4278 : SPIRV::OpConstantComposite;
4279 unsigned ContinuedOpc = HasSpecConst
4280 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
4281 : SPIRV::OpConstantCompositeContinuedINTEL;
4282 MachineIRBuilder MIR(
I);
4284 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
4286 for (
auto *Instr : Instructions) {
4287 Instr->setDebugLoc(
I.getDebugLoc());
4292 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4299 case Intrinsic::spv_assign_name: {
4300 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4301 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4302 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4303 i <
I.getNumExplicitOperands(); ++i) {
4304 MIB.
addImm(
I.getOperand(i).getImm());
4309 case Intrinsic::spv_switch: {
4310 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4311 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4312 if (
I.getOperand(i).isReg())
4313 MIB.
addReg(
I.getOperand(i).getReg());
4314 else if (
I.getOperand(i).isCImm())
4315 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4316 else if (
I.getOperand(i).isMBB())
4317 MIB.
addMBB(
I.getOperand(i).getMBB());
4324 case Intrinsic::spv_loop_merge: {
4325 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4326 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4327 if (
I.getOperand(i).isMBB())
4328 MIB.
addMBB(
I.getOperand(i).getMBB());
4335 case Intrinsic::spv_loop_control_intel: {
4337 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4338 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4343 case Intrinsic::spv_selection_merge: {
4345 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4346 assert(
I.getOperand(1).isMBB() &&
4347 "operand 1 to spv_selection_merge must be a basic block");
4348 MIB.
addMBB(
I.getOperand(1).getMBB());
4349 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4353 case Intrinsic::spv_cmpxchg:
4354 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4355 case Intrinsic::spv_unreachable:
4356 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4359 case Intrinsic::spv_alloca:
4360 return selectFrameIndex(ResVReg, ResType,
I);
4361 case Intrinsic::spv_alloca_array:
4362 return selectAllocaArray(ResVReg, ResType,
I);
4363 case Intrinsic::spv_assume:
4365 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4366 .
addUse(
I.getOperand(1).getReg())
4371 case Intrinsic::spv_expect:
4373 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4376 .
addUse(
I.getOperand(2).getReg())
4377 .
addUse(
I.getOperand(3).getReg())
4382 case Intrinsic::arithmetic_fence:
4383 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4384 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4387 .
addUse(
I.getOperand(2).getReg())
4391 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4393 case Intrinsic::spv_thread_id:
4399 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4401 case Intrinsic::spv_thread_id_in_group:
4407 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4409 case Intrinsic::spv_group_id:
4415 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4417 case Intrinsic::spv_flattened_thread_id_in_group:
4424 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4426 case Intrinsic::spv_workgroup_size:
4427 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4429 case Intrinsic::spv_global_size:
4430 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4432 case Intrinsic::spv_global_offset:
4433 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4435 case Intrinsic::spv_num_workgroups:
4436 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4438 case Intrinsic::spv_subgroup_size:
4439 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4441 case Intrinsic::spv_num_subgroups:
4442 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4444 case Intrinsic::spv_subgroup_id:
4445 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4446 case Intrinsic::spv_subgroup_local_invocation_id:
4447 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4448 ResVReg, ResType,
I);
4449 case Intrinsic::spv_subgroup_max_size:
4450 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4452 case Intrinsic::spv_fdot:
4453 return selectFloatDot(ResVReg, ResType,
I);
4454 case Intrinsic::spv_udot:
4455 case Intrinsic::spv_sdot:
4456 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4458 return selectIntegerDot(ResVReg, ResType,
I,
4459 IID == Intrinsic::spv_sdot);
4460 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4461 case Intrinsic::spv_dot4add_i8packed:
4462 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4464 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4465 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4466 case Intrinsic::spv_dot4add_u8packed:
4467 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4469 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4470 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4471 case Intrinsic::spv_all:
4472 return selectAll(ResVReg, ResType,
I);
4473 case Intrinsic::spv_any:
4474 return selectAny(ResVReg, ResType,
I);
4475 case Intrinsic::spv_cross:
4476 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4477 case Intrinsic::spv_distance:
4478 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4479 case Intrinsic::spv_lerp:
4480 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4481 case Intrinsic::spv_length:
4482 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4483 case Intrinsic::spv_degrees:
4484 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4485 case Intrinsic::spv_faceforward:
4486 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4487 case Intrinsic::spv_frac:
4488 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4489 case Intrinsic::spv_isinf:
4490 return selectOpIsInf(ResVReg, ResType,
I);
4491 case Intrinsic::spv_isnan:
4492 return selectOpIsNan(ResVReg, ResType,
I);
4493 case Intrinsic::spv_normalize:
4494 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4495 case Intrinsic::spv_refract:
4496 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4497 case Intrinsic::spv_reflect:
4498 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4499 case Intrinsic::spv_rsqrt:
4500 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4501 case Intrinsic::spv_sign:
4502 return selectSign(ResVReg, ResType,
I);
4503 case Intrinsic::spv_smoothstep:
4504 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4505 case Intrinsic::spv_firstbituhigh:
4506 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4507 case Intrinsic::spv_firstbitshigh:
4508 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4509 case Intrinsic::spv_firstbitlow:
4510 return selectFirstBitLow(ResVReg, ResType,
I);
4511 case Intrinsic::spv_group_memory_barrier:
4512 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4513 SPIRV::MemorySemantics::WorkgroupMemory,
4515 case Intrinsic::spv_group_memory_barrier_with_group_sync:
4516 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4517 SPIRV::MemorySemantics::WorkgroupMemory,
4519 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4520 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4521 SPIRV::StorageClass::StorageClass ResSC =
4525 "Generic storage class");
4526 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4534 case Intrinsic::spv_lifetime_start:
4535 case Intrinsic::spv_lifetime_end: {
4536 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
4537 : SPIRV::OpLifetimeStop;
4538 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
4539 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
4548 case Intrinsic::spv_saturate:
4549 return selectSaturate(ResVReg, ResType,
I);
4550 case Intrinsic::spv_nclamp:
4551 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
4552 case Intrinsic::spv_uclamp:
4553 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
4554 case Intrinsic::spv_sclamp:
4555 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
4556 case Intrinsic::spv_subgroup_prefix_bit_count:
4557 return selectWavePrefixBitCount(ResVReg, ResType,
I);
4558 case Intrinsic::spv_wave_active_countbits:
4559 return selectWaveActiveCountBits(ResVReg, ResType,
I);
4560 case Intrinsic::spv_wave_all_equal:
4561 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
4562 case Intrinsic::spv_wave_all:
4563 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
4564 case Intrinsic::spv_wave_any:
4565 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
4566 case Intrinsic::spv_subgroup_ballot:
4567 return selectWaveOpInst(ResVReg, ResType,
I,
4568 SPIRV::OpGroupNonUniformBallot);
4569 case Intrinsic::spv_wave_is_first_lane:
4570 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
4571 case Intrinsic::spv_wave_reduce_or:
4572 return selectWaveReduceOp(ResVReg, ResType,
I,
4573 SPIRV::OpGroupNonUniformBitwiseOr);
4574 case Intrinsic::spv_wave_reduce_xor:
4575 return selectWaveReduceOp(ResVReg, ResType,
I,
4576 SPIRV::OpGroupNonUniformBitwiseXor);
4577 case Intrinsic::spv_wave_reduce_and:
4578 return selectWaveReduceOp(ResVReg, ResType,
I,
4579 SPIRV::OpGroupNonUniformBitwiseAnd);
4580 case Intrinsic::spv_wave_reduce_umax:
4581 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
4582 case Intrinsic::spv_wave_reduce_max:
4583 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
4584 case Intrinsic::spv_wave_reduce_umin:
4585 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
4586 case Intrinsic::spv_wave_reduce_min:
4587 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
4588 case Intrinsic::spv_wave_reduce_sum:
4589 return selectWaveReduceSum(ResVReg, ResType,
I);
4590 case Intrinsic::spv_wave_product:
4591 return selectWaveReduceProduct(ResVReg, ResType,
I);
4592 case Intrinsic::spv_wave_readlane:
4593 return selectWaveOpInst(ResVReg, ResType,
I,
4594 SPIRV::OpGroupNonUniformShuffle);
4595 case Intrinsic::spv_wave_prefix_sum:
4596 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
4597 case Intrinsic::spv_wave_prefix_product:
4598 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
4599 case Intrinsic::spv_quad_read_across_x: {
4600 return selectQuadSwap(ResVReg, ResType,
I, 0);
4602 case Intrinsic::spv_quad_read_across_y: {
4603 return selectQuadSwap(ResVReg, ResType,
I, 1);
4605 case Intrinsic::spv_step:
4606 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
4607 case Intrinsic::spv_radians:
4608 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
4612 case Intrinsic::instrprof_increment:
4613 case Intrinsic::instrprof_increment_step:
4614 case Intrinsic::instrprof_value_profile:
4617 case Intrinsic::spv_value_md:
4619 case Intrinsic::spv_resource_handlefrombinding: {
4620 return selectHandleFromBinding(ResVReg, ResType,
I);
4622 case Intrinsic::spv_resource_counterhandlefrombinding:
4623 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
4624 case Intrinsic::spv_resource_updatecounter:
4625 return selectUpdateCounter(ResVReg, ResType,
I);
4626 case Intrinsic::spv_resource_store_typedbuffer: {
4627 return selectImageWriteIntrinsic(
I);
4629 case Intrinsic::spv_resource_load_typedbuffer: {
4630 return selectReadImageIntrinsic(ResVReg, ResType,
I);
4632 case Intrinsic::spv_resource_load_level: {
4633 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
4635 case Intrinsic::spv_resource_sample:
4636 case Intrinsic::spv_resource_sample_clamp:
4637 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
4638 case Intrinsic::spv_resource_samplebias:
4639 case Intrinsic::spv_resource_samplebias_clamp:
4640 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
4641 case Intrinsic::spv_resource_samplegrad:
4642 case Intrinsic::spv_resource_samplegrad_clamp:
4643 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
4644 case Intrinsic::spv_resource_samplelevel:
4645 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
4646 case Intrinsic::spv_resource_samplecmp:
4647 case Intrinsic::spv_resource_samplecmp_clamp:
4648 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
4649 case Intrinsic::spv_resource_samplecmplevelzero:
4650 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
4651 case Intrinsic::spv_resource_gather:
4652 case Intrinsic::spv_resource_gather_cmp:
4653 return selectGatherIntrinsic(ResVReg, ResType,
I);
4654 case Intrinsic::spv_resource_getpointer: {
4655 return selectResourceGetPointer(ResVReg, ResType,
I);
4657 case Intrinsic::spv_pushconstant_getpointer: {
4658 return selectPushConstantGetPointer(ResVReg, ResType,
I);
4660 case Intrinsic::spv_discard: {
4661 return selectDiscard(ResVReg, ResType,
I);
4663 case Intrinsic::spv_resource_nonuniformindex: {
4664 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
4666 case Intrinsic::spv_unpackhalf2x16: {
4667 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
4669 case Intrinsic::spv_packhalf2x16: {
4670 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
4672 case Intrinsic::spv_ddx:
4673 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
4674 case Intrinsic::spv_ddy:
4675 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
4676 case Intrinsic::spv_ddx_coarse:
4677 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
4678 case Intrinsic::spv_ddy_coarse:
4679 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
4680 case Intrinsic::spv_ddx_fine:
4681 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
4682 case Intrinsic::spv_ddy_fine:
4683 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
4684 case Intrinsic::spv_fwidth:
4685 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
4686 case Intrinsic::spv_masked_gather:
4687 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4688 return selectMaskedGather(ResVReg, ResType,
I);
4689 return diagnoseUnsupported(
4690 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
4691 case Intrinsic::spv_masked_scatter:
4692 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4693 return selectMaskedScatter(
I);
4694 return diagnoseUnsupported(
4695 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
4697 std::string DiagMsg;
4698 raw_string_ostream OS(DiagMsg);
4700 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
4707bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
4708 SPIRVTypeInst ResType,
4709 MachineInstr &
I)
const {
4712 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
4719bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
4720 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4722 assert(Intr.getIntrinsicID() ==
4723 Intrinsic::spv_resource_counterhandlefrombinding);
4726 Register MainHandleReg = Intr.getOperand(2).getReg();
4728 assert(MainHandleDef->getIntrinsicID() ==
4729 Intrinsic::spv_resource_handlefrombinding);
4733 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
4734 Register IndexReg = MainHandleDef->getOperand(5).getReg();
4735 std::string CounterName =
4740 MachineIRBuilder MIRBuilder(
I);
4742 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
4744 ArraySize, IndexReg, CounterName, MIRBuilder);
4746 return BuildCOPY(ResVReg, CounterVarReg,
I);
4749bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
4750 SPIRVTypeInst ResType,
4751 MachineInstr &
I)
const {
4753 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
4755 Register CounterHandleReg = Intr.getOperand(2).getReg();
4756 Register IncrReg = Intr.getOperand(3).getReg();
4763 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
4764 assert(CounterVarPointeeType &&
4765 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
4766 "Counter variable must be a struct");
4768 SPIRV::StorageClass::StorageBuffer &&
4769 "Counter variable must be in the storage buffer storage class");
4771 "Counter variable must have exactly 1 member in the struct");
4772 const SPIRVTypeInst MemberType =
4775 "Counter variable struct must have a single i32 member");
4779 MachineIRBuilder MIRBuilder(
I);
4781 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
4784 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
4790 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4793 .
addUse(CounterHandleReg)
4800 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
4803 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
4806 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
4815 return BuildCOPY(ResVReg, AtomicRes,
I);
4823 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
4831bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
4832 SPIRVTypeInst ResType,
4833 MachineInstr &
I)
const {
4841 Register ImageReg =
I.getOperand(2).getReg();
4849 Register IdxReg =
I.getOperand(3).getReg();
4851 MachineInstr &Pos =
I;
4853 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
4857bool SPIRVInstructionSelector::generateSampleImage(
4860 DebugLoc Loc, MachineInstr &Pos)
const {
4871 if (!loadHandleBeforePosition(NewSamplerReg,
4877 MachineIRBuilder MIRBuilder(Pos);
4890 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
4891 ImOps.Lod.has_value();
4892 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
4893 : SPIRV::OpImageSampleImplicitLod;
4895 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
4896 : SPIRV::OpImageSampleDrefImplicitLod;
4905 MIB.
addUse(*ImOps.Compare);
4907 uint32_t ImageOperands = 0;
4909 ImageOperands |= SPIRV::ImageOperand::Bias;
4911 ImageOperands |= SPIRV::ImageOperand::Lod;
4912 if (ImOps.GradX && ImOps.GradY)
4913 ImageOperands |= SPIRV::ImageOperand::Grad;
4914 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
4916 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4919 "Non-constant offsets are not supported in sample instructions.");
4923 ImageOperands |= SPIRV::ImageOperand::MinLod;
4925 if (ImageOperands != 0) {
4926 MIB.
addImm(ImageOperands);
4927 if (ImageOperands & SPIRV::ImageOperand::Bias)
4929 if (ImageOperands & SPIRV::ImageOperand::Lod)
4931 if (ImageOperands & SPIRV::ImageOperand::Grad) {
4932 MIB.
addUse(*ImOps.GradX);
4933 MIB.
addUse(*ImOps.GradY);
4936 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4937 MIB.
addUse(*ImOps.Offset);
4938 if (ImageOperands & SPIRV::ImageOperand::MinLod)
4939 MIB.
addUse(*ImOps.MinLod);
4946bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
4947 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4948 Register ImageReg =
I.getOperand(2).getReg();
4949 Register SamplerReg =
I.getOperand(3).getReg();
4950 Register CoordinateReg =
I.getOperand(4).getReg();
4951 ImageOperands ImOps;
4952 if (
I.getNumOperands() > 5)
4953 ImOps.Offset =
I.getOperand(5).getReg();
4954 if (
I.getNumOperands() > 6)
4955 ImOps.MinLod =
I.getOperand(6).getReg();
4956 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4957 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4960bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
4961 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4962 Register ImageReg =
I.getOperand(2).getReg();
4963 Register SamplerReg =
I.getOperand(3).getReg();
4964 Register CoordinateReg =
I.getOperand(4).getReg();
4965 ImageOperands ImOps;
4966 ImOps.Bias =
I.getOperand(5).getReg();
4967 if (
I.getNumOperands() > 6)
4968 ImOps.Offset =
I.getOperand(6).getReg();
4969 if (
I.getNumOperands() > 7)
4970 ImOps.MinLod =
I.getOperand(7).getReg();
4971 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4972 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4975bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
4976 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4977 Register ImageReg =
I.getOperand(2).getReg();
4978 Register SamplerReg =
I.getOperand(3).getReg();
4979 Register CoordinateReg =
I.getOperand(4).getReg();
4980 ImageOperands ImOps;
4981 ImOps.GradX =
I.getOperand(5).getReg();
4982 ImOps.GradY =
I.getOperand(6).getReg();
4983 if (
I.getNumOperands() > 7)
4984 ImOps.Offset =
I.getOperand(7).getReg();
4985 if (
I.getNumOperands() > 8)
4986 ImOps.MinLod =
I.getOperand(8).getReg();
4987 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4988 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4991bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
4992 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4993 Register ImageReg =
I.getOperand(2).getReg();
4994 Register SamplerReg =
I.getOperand(3).getReg();
4995 Register CoordinateReg =
I.getOperand(4).getReg();
4996 ImageOperands ImOps;
4997 ImOps.Lod =
I.getOperand(5).getReg();
4998 if (
I.getNumOperands() > 6)
4999 ImOps.Offset =
I.getOperand(6).getReg();
5000 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5001 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5004bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
5005 SPIRVTypeInst ResType,
5006 MachineInstr &
I)
const {
5007 Register ImageReg =
I.getOperand(2).getReg();
5008 Register SamplerReg =
I.getOperand(3).getReg();
5009 Register CoordinateReg =
I.getOperand(4).getReg();
5010 ImageOperands ImOps;
5011 ImOps.Compare =
I.getOperand(5).getReg();
5012 if (
I.getNumOperands() > 6)
5013 ImOps.Offset =
I.getOperand(6).getReg();
5014 if (
I.getNumOperands() > 7)
5015 ImOps.MinLod =
I.getOperand(7).getReg();
5016 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5017 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5020bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
5021 SPIRVTypeInst ResType,
5022 MachineInstr &
I)
const {
5023 Register ImageReg =
I.getOperand(2).getReg();
5024 Register CoordinateReg =
I.getOperand(3).getReg();
5025 Register LodReg =
I.getOperand(4).getReg();
5027 ImageOperands ImOps;
5029 if (
I.getNumOperands() > 5)
5030 ImOps.Offset =
I.getOperand(5).getReg();
5042 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
5043 I.getDebugLoc(),
I, &ImOps);
5046bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
5047 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5048 Register ImageReg =
I.getOperand(2).getReg();
5049 Register SamplerReg =
I.getOperand(3).getReg();
5050 Register CoordinateReg =
I.getOperand(4).getReg();
5051 ImageOperands ImOps;
5052 ImOps.Compare =
I.getOperand(5).getReg();
5053 if (
I.getNumOperands() > 6)
5054 ImOps.Offset =
I.getOperand(6).getReg();
5057 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5058 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5061bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
5062 SPIRVTypeInst ResType,
5063 MachineInstr &
I)
const {
5064 Register ImageReg =
I.getOperand(2).getReg();
5065 Register SamplerReg =
I.getOperand(3).getReg();
5066 Register CoordinateReg =
I.getOperand(4).getReg();
5069 "ImageReg is not an image type.");
5074 ComponentOrCompareReg =
I.getOperand(5).getReg();
5075 OffsetReg =
I.getOperand(6).getReg();
5078 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
5082 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5083 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
5084 Dim != SPIRV::Dim::DIM_Rect) {
5086 "Gather operations are only supported for 2D, Cube, and Rect images.");
5093 if (!loadHandleBeforePosition(
5098 MachineIRBuilder MIRBuilder(
I);
5099 SPIRVTypeInst SampledImageType =
5104 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5112 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
5114 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
5116 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5121 .
addUse(ComponentOrCompareReg);
5123 uint32_t ImageOperands = 0;
5124 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
5125 if (Dim == SPIRV::Dim::DIM_Cube) {
5127 "Gather operations with offset are not supported for Cube images.");
5131 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5133 ImageOperands |= SPIRV::ImageOperand::Offset;
5137 if (ImageOperands != 0) {
5138 MIB.
addImm(ImageOperands);
5140 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5148bool SPIRVInstructionSelector::generateImageReadOrFetch(
5151 const ImageOperands *ImOps)
const {
5154 "ImageReg is not an image type.");
5156 bool IsSignedInteger =
5161 bool IsFetch = (SampledOp.getImm() == 1);
5163 auto AddOperands = [&](MachineInstrBuilder &MIB) {
5164 uint32_t ImageOperandsMask = 0;
5165 if (IsSignedInteger)
5166 ImageOperandsMask |= 0x1000;
5168 if (IsFetch && ImOps) {
5170 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
5171 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
5173 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
5175 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5179 if (ImageOperandsMask != 0) {
5180 MIB.
addImm(ImageOperandsMask);
5181 if (IsFetch && ImOps) {
5184 if (ImOps->Offset &&
5185 (ImageOperandsMask &
5186 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5187 MIB.
addUse(*ImOps->Offset);
5193 if (ResultSize == 4) {
5196 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5203 BMI.constrainAllUses(
TII,
TRI, RBI);
5207 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5211 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5217 BMI.constrainAllUses(
TII,
TRI, RBI);
5219 if (ResultSize == 1) {
5228 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
5231bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
5232 SPIRVTypeInst ResType,
5233 MachineInstr &
I)
const {
5234 Register ResourcePtr =
I.getOperand(2).getReg();
5236 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
5245 MachineIRBuilder MIRBuilder(
I);
5247 Register IndexReg =
I.getOperand(3).getReg();
5250 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5260bool SPIRVInstructionSelector::selectPushConstantGetPointer(
5261 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5266bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
5267 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5268 Register ObjReg =
I.getOperand(2).getReg();
5269 if (!BuildCOPY(ResVReg, ObjReg,
I))
5279 decorateUsesAsNonUniform(ResVReg);
5283void SPIRVInstructionSelector::decorateUsesAsNonUniform(
5286 while (WorkList.
size() > 0) {
5290 bool IsDecorated =
false;
5292 if (
Use.getOpcode() == SPIRV::OpDecorate &&
5293 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
5299 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
5301 if (ResultReg == CurrentReg)
5309 SPIRV::Decoration::NonUniformEXT, {});
5314bool SPIRVInstructionSelector::extractSubvector(
5316 MachineInstr &InsertionPoint)
const {
5318 [[maybe_unused]] uint64_t InputSize =
5321 assert(InputSize > 1 &&
"The input must be a vector.");
5322 assert(ResultSize > 1 &&
"The result must be a vector.");
5323 assert(ResultSize < InputSize &&
5324 "Cannot extract more element than there are in the input.");
5327 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
5328 for (uint64_t
I = 0;
I < ResultSize;
I++) {
5331 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5340 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
5342 TII.get(SPIRV::OpCompositeConstruct))
5346 for (
Register ComponentReg : ComponentRegisters)
5347 MIB.
addUse(ComponentReg);
5352bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
5353 MachineInstr &
I)
const {
5360 Register ImageReg =
I.getOperand(1).getReg();
5368 Register CoordinateReg =
I.getOperand(2).getReg();
5369 Register DataReg =
I.getOperand(3).getReg();
5372 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
5380Register SPIRVInstructionSelector::buildPointerToResource(
5381 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
5382 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
5383 StringRef Name, MachineIRBuilder MIRBuilder)
const {
5385 if (ArraySize == 1) {
5386 SPIRVTypeInst PtrType =
5389 "SpirvResType did not have an explicit layout.");
5394 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
5395 SPIRVTypeInst VarPointerType =
5398 VarPointerType, Set,
Binding, Name, MIRBuilder);
5400 SPIRVTypeInst ResPointerType =
5413bool SPIRVInstructionSelector::selectFirstBitSet16(
5414 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5415 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
5417 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
5421 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
5424bool SPIRVInstructionSelector::selectFirstBitSet32(
5426 unsigned BitSetOpcode)
const {
5427 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5430 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5437bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
5439 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5446 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
5448 MachineIRBuilder MIRBuilder(
I);
5451 SPIRVTypeInst I64x2Type =
5453 SPIRVTypeInst Vec2ResType =
5456 std::vector<Register> PartialRegs;
5459 unsigned CurrentComponent = 0;
5460 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
5466 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5467 TII.get(SPIRV::OpVectorShuffle))
5472 .
addImm(CurrentComponent)
5473 .
addImm(CurrentComponent + 1);
5480 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType,
I, BitSetResult,
5481 BitSetOpcode, SwapPrimarySide))
5484 PartialRegs.push_back(SubVecBitSetReg);
5488 if (CurrentComponent != ComponentCount) {
5494 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
5495 SPIRV::OpVectorExtractDynamic))
5501 if (!selectFirstBitSet64(FinalElemBitSetReg,
BaseType,
I, FinalElemReg,
5502 BitSetOpcode, SwapPrimarySide))
5505 PartialRegs.push_back(FinalElemBitSetReg);
5510 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
5511 SPIRV::OpCompositeConstruct);
5514bool SPIRVInstructionSelector::selectFirstBitSet64(
5516 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5529 if (ComponentCount > 2) {
5530 return selectFirstBitSet64Overflow(ResVReg, ResType,
I, SrcReg,
5531 BitSetOpcode, SwapPrimarySide);
5535 MachineIRBuilder MIRBuilder(
I);
5537 BaseType, 2 * ComponentCount, MIRBuilder,
false);
5541 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
5547 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
5554 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
5557 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
5558 SPIRV::OpVectorExtractDynamic))
5560 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
5561 SPIRV::OpVectorExtractDynamic))
5565 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5566 TII.get(SPIRV::OpVectorShuffle))
5574 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
5580 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5581 TII.get(SPIRV::OpVectorShuffle))
5589 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
5609 SelectOp = SPIRV::OpSelectSISCond;
5610 AddOp = SPIRV::OpIAddS;
5618 SelectOp = SPIRV::OpSelectVIVCond;
5619 AddOp = SPIRV::OpIAddV;
5625 Register RegSecondaryOffset = Reg0;
5629 if (SwapPrimarySide) {
5630 PrimaryReg = LowReg;
5631 SecondaryReg = HighReg;
5632 RegPrimaryOffset = Reg0;
5633 RegSecondaryOffset = Reg32;
5638 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
5639 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
5644 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
5645 SPIRV::OpINotEqual))
5652 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
5653 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
5658 if (SwapPrimarySide) {
5660 if (!selectOpWithSrcs(RegAdd, ResType,
I,
5661 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
5672 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
5673 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
5678 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
5679 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
5682 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
5686bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
5687 SPIRVTypeInst ResType,
5689 bool IsSigned)
const {
5691 Register OpReg =
I.getOperand(2).getReg();
5694 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5695 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
5699 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5701 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5703 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5707 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
5711bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
5712 SPIRVTypeInst ResType,
5713 MachineInstr &
I)
const {
5715 Register OpReg =
I.getOperand(2).getReg();
5720 unsigned ExtendOpcode = SPIRV::OpUConvert;
5721 unsigned BitSetOpcode = GL::FindILsb;
5725 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5727 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5729 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5736bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
5737 SPIRVTypeInst ResType,
5738 MachineInstr &
I)
const {
5742 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
5745 .
addUse(
I.getOperand(2).getReg())
5748 unsigned Alignment =
I.getOperand(3).getImm();
5754bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
5755 SPIRVTypeInst ResType,
5756 MachineInstr &
I)
const {
5760 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
5763 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
5766 unsigned Alignment =
I.getOperand(2).getImm();
5773bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
5778 const MachineInstr *PrevI =
I.getPrevNode();
5780 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
5784 .
addMBB(
I.getOperand(0).getMBB())
5789 .
addMBB(
I.getOperand(0).getMBB())
5794bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
5805 const MachineInstr *NextI =
I.getNextNode();
5807 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
5813 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
5815 .
addUse(
I.getOperand(0).getReg())
5816 .
addMBB(
I.getOperand(1).getMBB())
5822bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
5823 MachineInstr &
I)
const {
5825 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
5827 const unsigned NumOps =
I.getNumOperands();
5828 for (
unsigned i = 1; i <
NumOps; i += 2) {
5829 MIB.
addUse(
I.getOperand(i + 0).getReg());
5830 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
5836bool SPIRVInstructionSelector::selectGlobalValue(
5837 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
5839 MachineIRBuilder MIRBuilder(
I);
5840 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
5843 std::string GlobalIdent;
5845 unsigned &
ID = UnnamedGlobalIDs[GV];
5847 ID = UnnamedGlobalIDs.
size();
5848 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
5874 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
5881 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
5886 MachineInstrBuilder MIB1 =
5887 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5890 MachineInstrBuilder MIB2 =
5892 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
5896 GR.
add(ConstVal, MIB2);
5904 MachineInstrBuilder MIB3 =
5905 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5908 GR.
add(ConstVal, MIB3);
5912 assert(NewReg != ResVReg);
5913 return BuildCOPY(ResVReg, NewReg,
I);
5923 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
5929 SPIRVTypeInst ResType =
5933 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
5938 if (
GlobalVar->isExternallyInitialized() &&
5939 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
5940 constexpr unsigned ReadWriteINTEL = 3u;
5943 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
5949bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
5950 SPIRVTypeInst ResType,
5951 MachineInstr &
I)
const {
5953 return selectExtInst(ResVReg, ResType,
I, CL::log10);
5961 MachineIRBuilder MIRBuilder(
I);
5966 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5969 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5971 .
add(
I.getOperand(1))
5976 ResType->
getOpcode() == SPIRV::OpTypeFloat);
5978 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
5986 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
5987 ? SPIRV::OpVectorTimesScalar
5998bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
5999 SPIRVTypeInst ResType,
6000 MachineInstr &
I)
const {
6003 return selectExtInst(ResVReg, ResType,
I, CL::pown);
6009 Register ExpReg =
I.getOperand(2).getReg();
6011 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
6012 SPIRV::OpConvertSToF))
6014 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
6021bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
6022 SPIRVTypeInst ResType,
6023 MachineInstr &
I)
const {
6039 MachineIRBuilder MIRBuilder(
I);
6042 ResType, MIRBuilder, SPIRV::StorageClass::Function);
6054 MachineBasicBlock &EntryBB =
I.getMF()->front();
6058 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
6061 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
6067 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6070 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
6073 .
add(
I.getOperand(
I.getNumExplicitDefs()))
6077 Register IntegralPartReg =
I.getOperand(1).getReg();
6078 if (IntegralPartReg.
isValid()) {
6080 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6091 assert(
false &&
"GLSL::Modf is deprecated.");
6102bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
6103 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6104 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6105 MachineIRBuilder MIRBuilder(
I);
6106 const SPIRVTypeInst Vec3Ty =
6109 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
6121 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6125 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6131 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6138 assert(
I.getOperand(2).isReg());
6139 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
6143 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6154bool SPIRVInstructionSelector::loadBuiltinInputID(
6155 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6156 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6157 MachineIRBuilder MIRBuilder(
I);
6159 ResType, MIRBuilder, SPIRV::StorageClass::Input);
6174 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6178 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6187SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
6188 MachineInstr &
I)
const {
6189 MachineIRBuilder MIRBuilder(
I);
6190 if (
Type->getOpcode() != SPIRV::OpTypeVector)
6194 if (VectorSize == 4)
6202bool SPIRVInstructionSelector::loadHandleBeforePosition(
6203 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
6204 MachineInstr &Pos)
const {
6207 Intrinsic::spv_resource_handlefrombinding);
6215 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
6216 MachineIRBuilder MIRBuilder(HandleDef);
6217 SPIRVTypeInst VarType = ResType;
6218 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
6220 if (IsStructuredBuffer) {
6225 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
6227 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
6230 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
6231 ArraySize, IndexReg, Name, MIRBuilder);
6235 uint32_t LoadOpcode =
6236 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
6246void SPIRVInstructionSelector::errorIfInstrOutsideShader(
6247 MachineInstr &
I)
const {
6249 std::string DiagMsg;
6250 raw_string_ostream OS(DiagMsg);
6251 I.print(OS,
true,
false,
false,
false);
6252 DiagMsg +=
" is only supported in shaders.\n";
6258InstructionSelector *
6262 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...