33#include "llvm/IR/IntrinsicsSPIRV.h"
37#define DEBUG_TYPE "spirv-isel"
44 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
49 std::optional<Register> Bias;
50 std::optional<Register>
Offset;
51 std::optional<Register> MinLod;
52 std::optional<Register> GradX;
53 std::optional<Register> GradY;
54 std::optional<Register> Lod;
55 std::optional<Register> Compare;
58llvm::SPIRV::SelectionControl::SelectionControl
59getSelectionOperandForImm(
int Imm) {
61 return SPIRV::SelectionControl::Flatten;
63 return SPIRV::SelectionControl::DontFlatten;
65 return SPIRV::SelectionControl::None;
69#define GET_GLOBALISEL_PREDICATE_BITSET
70#include "SPIRVGenGlobalISel.inc"
71#undef GET_GLOBALISEL_PREDICATE_BITSET
98#define GET_GLOBALISEL_PREDICATES_DECL
99#include "SPIRVGenGlobalISel.inc"
100#undef GET_GLOBALISEL_PREDICATES_DECL
102#define GET_GLOBALISEL_TEMPORARIES_DECL
103#include "SPIRVGenGlobalISel.inc"
104#undef GET_GLOBALISEL_TEMPORARIES_DECL
128 unsigned BitSetOpcode)
const;
132 unsigned BitSetOpcode)
const;
136 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
140 unsigned BitSetOpcode,
141 bool SwapPrimarySide)
const;
148 unsigned Opcode)
const;
151 unsigned Opcode)
const;
170 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
181 unsigned OpType)
const;
236 template <
bool Signed>
239 template <
bool Signed>
246 template <
typename PickOpcodeFn>
249 PickOpcodeFn &&PickOpcode)
const;
266 template <
typename PickOpcodeFn>
269 PickOpcodeFn &&PickOpcode)
const;
287 bool IsSigned)
const;
289 bool IsSigned,
unsigned Opcode)
const;
291 bool IsSigned)
const;
297 bool IsSigned)
const;
336 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
337 bool useMISrc =
true,
339 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
340 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
341 bool useMISrc =
true,
343 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
344 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
345 bool setMIFlags =
true,
bool useMISrc =
true,
347 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
348 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
349 bool useMISrc =
true,
352 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
353 MachineInstr &
I)
const;
355 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
356 MachineInstr &
I)
const;
358 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
359 MachineInstr &
I)
const;
361 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
362 MachineInstr &
I,
unsigned Opcode)
const;
364 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
365 MachineInstr &
I)
const;
367 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
368 MachineInstr &
I)
const;
372 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
373 MachineInstr &
I)
const;
375 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
376 MachineInstr &
I)
const;
378 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
379 MachineInstr &
I)
const;
380 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
381 MachineInstr &
I)
const;
382 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
383 MachineInstr &
I)
const;
384 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
385 MachineInstr &
I)
const;
386 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
387 MachineInstr &
I)
const;
388 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
389 MachineInstr &
I)
const;
390 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
391 MachineInstr &
I)
const;
392 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
393 SPIRVTypeInst ResType,
394 MachineInstr &
I)
const;
395 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
396 MachineInstr &
I)
const;
397 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
398 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
399 MachineInstr &
I)
const;
400 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
401 MachineInstr &
I)
const;
402 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
403 MachineInstr &
I)
const;
404 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
405 MachineInstr &
I)
const;
406 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
407 MachineInstr &
I)
const;
408 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
409 MachineInstr &
I)
const;
410 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
411 MachineInstr &
I)
const;
412 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
414 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
415 MachineInstr &
I,
const unsigned DPdOpCode)
const;
417 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
418 SPIRVTypeInst ResType =
nullptr)
const;
420 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
421 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
422 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
424 MachineInstr &
I)
const;
425 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
427 bool wrapIntoSpecConstantOp(MachineInstr &
I,
430 Register getUcharPtrTypeReg(MachineInstr &
I,
431 SPIRV::StorageClass::StorageClass SC)
const;
432 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
434 uint32_t Opcode)
const;
435 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
436 SPIRVTypeInst SrcPtrTy)
const;
437 Register buildPointerToResource(SPIRVTypeInst ResType,
438 SPIRV::StorageClass::StorageClass SC,
439 uint32_t Set, uint32_t
Binding,
440 uint32_t ArraySize,
Register IndexReg,
442 MachineIRBuilder MIRBuilder)
const;
443 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
444 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
445 Register &ReadReg, MachineInstr &InsertionPoint)
const;
446 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
449 const ImageOperands *ImOps =
nullptr)
const;
450 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
452 Register CoordinateReg,
const ImageOperands &ImOps,
455 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
456 Register ResVReg, SPIRVTypeInst ResType,
457 MachineInstr &
I)
const;
458 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
459 Register ResVReg, SPIRVTypeInst ResType,
460 MachineInstr &
I)
const;
461 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
462 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
463 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
464 void errorIfInstrOutsideShader(MachineInstr &
I)
const;
467bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
469 if (
TET->getTargetExtName() ==
"spirv.Image") {
472 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
473 return TET->getTypeParameter(0)->isIntegerTy();
477#define GET_GLOBALISEL_IMPL
478#include "SPIRVGenGlobalISel.inc"
479#undef GET_GLOBALISEL_IMPL
485 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
488#include
"SPIRVGenGlobalISel.inc"
491#include
"SPIRVGenGlobalISel.inc"
503 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
507void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
508 if (HasVRegsReset == &MF)
523 for (
const auto &
MBB : MF) {
524 for (
const auto &
MI :
MBB) {
527 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
531 LLT DstType = MRI.
getType(DstReg);
533 LLT SrcType = MRI.
getType(SrcReg);
534 if (DstType != SrcType)
539 if (DstRC != SrcRC && SrcRC)
551 while (!Stack.empty()) {
556 switch (
MI->getOpcode()) {
557 case TargetOpcode::G_INTRINSIC:
558 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
559 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
562 if (IntrID != Intrinsic::spv_const_composite &&
563 IntrID != Intrinsic::spv_undef)
567 case TargetOpcode::G_BUILD_VECTOR:
568 case TargetOpcode::G_SPLAT_VECTOR:
570 i < OpDef->getNumOperands(); i++) {
575 Stack.push_back(OpNestedDef);
578 case TargetOpcode::G_CONSTANT:
579 case TargetOpcode::G_FCONSTANT:
580 case TargetOpcode::G_IMPLICIT_DEF:
581 case SPIRV::OpConstantTrue:
582 case SPIRV::OpConstantFalse:
583 case SPIRV::OpConstantI:
584 case SPIRV::OpConstantF:
585 case SPIRV::OpConstantComposite:
586 case SPIRV::OpConstantCompositeContinuedINTEL:
587 case SPIRV::OpConstantSampler:
588 case SPIRV::OpConstantNull:
590 case SPIRV::OpConstantFunctionPointerINTEL:
617 case Intrinsic::spv_all:
618 case Intrinsic::spv_alloca:
619 case Intrinsic::spv_any:
620 case Intrinsic::spv_bitcast:
621 case Intrinsic::spv_const_composite:
622 case Intrinsic::spv_cross:
623 case Intrinsic::spv_degrees:
624 case Intrinsic::spv_distance:
625 case Intrinsic::spv_extractelt:
626 case Intrinsic::spv_extractv:
627 case Intrinsic::spv_faceforward:
628 case Intrinsic::spv_fdot:
629 case Intrinsic::spv_firstbitlow:
630 case Intrinsic::spv_firstbitshigh:
631 case Intrinsic::spv_firstbituhigh:
632 case Intrinsic::spv_frac:
633 case Intrinsic::spv_gep:
634 case Intrinsic::spv_global_offset:
635 case Intrinsic::spv_global_size:
636 case Intrinsic::spv_group_id:
637 case Intrinsic::spv_insertelt:
638 case Intrinsic::spv_insertv:
639 case Intrinsic::spv_isinf:
640 case Intrinsic::spv_isnan:
641 case Intrinsic::spv_lerp:
642 case Intrinsic::spv_length:
643 case Intrinsic::spv_normalize:
644 case Intrinsic::spv_num_subgroups:
645 case Intrinsic::spv_num_workgroups:
646 case Intrinsic::spv_ptrcast:
647 case Intrinsic::spv_radians:
648 case Intrinsic::spv_reflect:
649 case Intrinsic::spv_refract:
650 case Intrinsic::spv_resource_getpointer:
651 case Intrinsic::spv_resource_handlefrombinding:
652 case Intrinsic::spv_resource_handlefromimplicitbinding:
653 case Intrinsic::spv_resource_nonuniformindex:
654 case Intrinsic::spv_resource_sample:
655 case Intrinsic::spv_rsqrt:
656 case Intrinsic::spv_saturate:
657 case Intrinsic::spv_sdot:
658 case Intrinsic::spv_sign:
659 case Intrinsic::spv_smoothstep:
660 case Intrinsic::spv_step:
661 case Intrinsic::spv_subgroup_id:
662 case Intrinsic::spv_subgroup_local_invocation_id:
663 case Intrinsic::spv_subgroup_max_size:
664 case Intrinsic::spv_subgroup_size:
665 case Intrinsic::spv_thread_id:
666 case Intrinsic::spv_thread_id_in_group:
667 case Intrinsic::spv_udot:
668 case Intrinsic::spv_undef:
669 case Intrinsic::spv_value_md:
670 case Intrinsic::spv_workgroup_size:
682 case SPIRV::OpTypeVoid:
683 case SPIRV::OpTypeBool:
684 case SPIRV::OpTypeInt:
685 case SPIRV::OpTypeFloat:
686 case SPIRV::OpTypeVector:
687 case SPIRV::OpTypeMatrix:
688 case SPIRV::OpTypeImage:
689 case SPIRV::OpTypeSampler:
690 case SPIRV::OpTypeSampledImage:
691 case SPIRV::OpTypeArray:
692 case SPIRV::OpTypeRuntimeArray:
693 case SPIRV::OpTypeStruct:
694 case SPIRV::OpTypeOpaque:
695 case SPIRV::OpTypePointer:
696 case SPIRV::OpTypeFunction:
697 case SPIRV::OpTypeEvent:
698 case SPIRV::OpTypeDeviceEvent:
699 case SPIRV::OpTypeReserveId:
700 case SPIRV::OpTypeQueue:
701 case SPIRV::OpTypePipe:
702 case SPIRV::OpTypeForwardPointer:
703 case SPIRV::OpTypePipeStorage:
704 case SPIRV::OpTypeNamedBarrier:
705 case SPIRV::OpTypeAccelerationStructureNV:
706 case SPIRV::OpTypeCooperativeMatrixNV:
707 case SPIRV::OpTypeCooperativeMatrixKHR:
717 if (
MI.getNumDefs() == 0)
720 for (
const auto &MO :
MI.all_defs()) {
722 if (
Reg.isPhysical()) {
727 if (
UseMI.getOpcode() != SPIRV::OpName) {
734 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
735 MI.isLifetimeMarker()) {
738 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
749 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
750 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
753 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
758 if (
MI.mayStore() ||
MI.isCall() ||
759 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
760 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
761 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
772 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
779void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
781 for (
const auto &MO :
MI.all_defs()) {
785 SmallVector<MachineInstr *, 4> UselessOpNames;
788 "There is still a use of the dead function.");
791 for (MachineInstr *OpNameMI : UselessOpNames) {
793 OpNameMI->eraseFromParent();
798void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
801 removeOpNamesForDeadMI(
MI);
802 MI.eraseFromParent();
805bool SPIRVInstructionSelector::select(MachineInstr &
I) {
806 resetVRegsType(*
I.getParent()->getParent());
808 assert(
I.getParent() &&
"Instruction should be in a basic block!");
809 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
814 removeDeadInstruction(
I);
821 if (Opcode == SPIRV::ASSIGN_TYPE) {
822 Register DstReg =
I.getOperand(0).getReg();
823 Register SrcReg =
I.getOperand(1).getReg();
826 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
827 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
828 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
829 Register SelectDstReg =
Def->getOperand(0).getReg();
830 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
832 assert(SuccessToSelectSelect);
834 Def->eraseFromParent();
841 bool Res = selectImpl(
I, *CoverageInfo);
843 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
844 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
848 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
860 }
else if (
I.getNumDefs() == 1) {
872 removeDeadInstruction(
I);
877 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
878 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
884 bool HasDefs =
I.getNumDefs() > 0;
887 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
888 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
889 if (spvSelect(ResVReg, ResType,
I)) {
891 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
902 case TargetOpcode::G_CONSTANT:
903 case TargetOpcode::G_FCONSTANT:
905 case TargetOpcode::G_SADDO:
906 case TargetOpcode::G_SSUBO:
913 MachineInstr &
I)
const {
916 if (DstRC != SrcRC && SrcRC)
918 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
925bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
926 SPIRVTypeInst ResType,
927 MachineInstr &
I)
const {
928 const unsigned Opcode =
I.getOpcode();
930 return selectImpl(
I, *CoverageInfo);
932 case TargetOpcode::G_CONSTANT:
933 case TargetOpcode::G_FCONSTANT:
934 return selectConst(ResVReg, ResType,
I);
935 case TargetOpcode::G_GLOBAL_VALUE:
936 return selectGlobalValue(ResVReg,
I);
937 case TargetOpcode::G_IMPLICIT_DEF:
938 return selectOpUndef(ResVReg, ResType,
I);
939 case TargetOpcode::G_FREEZE:
940 return selectFreeze(ResVReg, ResType,
I);
942 case TargetOpcode::G_INTRINSIC:
943 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
944 case TargetOpcode::G_INTRINSIC_CONVERGENT:
945 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
946 return selectIntrinsic(ResVReg, ResType,
I);
947 case TargetOpcode::G_BITREVERSE:
948 return selectBitreverse(ResVReg, ResType,
I);
950 case TargetOpcode::G_BUILD_VECTOR:
951 return selectBuildVector(ResVReg, ResType,
I);
952 case TargetOpcode::G_SPLAT_VECTOR:
953 return selectSplatVector(ResVReg, ResType,
I);
955 case TargetOpcode::G_SHUFFLE_VECTOR: {
956 MachineBasicBlock &BB = *
I.getParent();
957 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
960 .
addUse(
I.getOperand(1).getReg())
961 .
addUse(
I.getOperand(2).getReg());
962 for (
auto V :
I.getOperand(3).getShuffleMask())
967 case TargetOpcode::G_MEMMOVE:
968 case TargetOpcode::G_MEMCPY:
969 case TargetOpcode::G_MEMSET:
970 return selectMemOperation(ResVReg,
I);
972 case TargetOpcode::G_ICMP:
973 return selectICmp(ResVReg, ResType,
I);
974 case TargetOpcode::G_FCMP:
975 return selectFCmp(ResVReg, ResType,
I);
977 case TargetOpcode::G_FRAME_INDEX:
978 return selectFrameIndex(ResVReg, ResType,
I);
980 case TargetOpcode::G_LOAD:
981 return selectLoad(ResVReg, ResType,
I);
982 case TargetOpcode::G_STORE:
983 return selectStore(
I);
985 case TargetOpcode::G_BR:
986 return selectBranch(
I);
987 case TargetOpcode::G_BRCOND:
988 return selectBranchCond(
I);
990 case TargetOpcode::G_PHI:
991 return selectPhi(ResVReg,
I);
993 case TargetOpcode::G_FPTOSI:
994 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
995 case TargetOpcode::G_FPTOUI:
996 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
998 case TargetOpcode::G_FPTOSI_SAT:
999 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1000 case TargetOpcode::G_FPTOUI_SAT:
1001 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1003 case TargetOpcode::G_SITOFP:
1004 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1005 case TargetOpcode::G_UITOFP:
1006 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1008 case TargetOpcode::G_CTPOP:
1009 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitCount);
1010 case TargetOpcode::G_SMIN:
1011 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1012 case TargetOpcode::G_UMIN:
1013 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1015 case TargetOpcode::G_SMAX:
1016 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1017 case TargetOpcode::G_UMAX:
1018 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1020 case TargetOpcode::G_SCMP:
1021 return selectSUCmp(ResVReg, ResType,
I,
true);
1022 case TargetOpcode::G_UCMP:
1023 return selectSUCmp(ResVReg, ResType,
I,
false);
1024 case TargetOpcode::G_LROUND:
1025 case TargetOpcode::G_LLROUND: {
1028 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1030 regForLround, *(
I.getParent()->getParent()));
1032 CL::round, GL::Round,
false);
1034 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1041 case TargetOpcode::G_STRICT_FMA:
1042 case TargetOpcode::G_FMA: {
1045 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1048 .
addUse(
I.getOperand(1).getReg())
1049 .
addUse(
I.getOperand(2).getReg())
1050 .
addUse(
I.getOperand(3).getReg())
1055 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1058 case TargetOpcode::G_STRICT_FLDEXP:
1059 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1061 case TargetOpcode::G_FPOW:
1062 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1063 case TargetOpcode::G_FPOWI:
1064 return selectFpowi(ResVReg, ResType,
I);
1066 case TargetOpcode::G_FEXP:
1067 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1068 case TargetOpcode::G_FEXP2:
1069 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1070 case TargetOpcode::G_FEXP10:
1071 return selectExp10(ResVReg, ResType,
I);
1073 case TargetOpcode::G_FMODF:
1074 return selectModf(ResVReg, ResType,
I);
1075 case TargetOpcode::G_FSINCOS:
1076 return selectSincos(ResVReg, ResType,
I);
1078 case TargetOpcode::G_FLOG:
1079 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1080 case TargetOpcode::G_FLOG2:
1081 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1082 case TargetOpcode::G_FLOG10:
1083 return selectLog10(ResVReg, ResType,
I);
1085 case TargetOpcode::G_FABS:
1086 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1087 case TargetOpcode::G_ABS:
1088 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1090 case TargetOpcode::G_FMINNUM:
1091 case TargetOpcode::G_FMINIMUM:
1092 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1093 case TargetOpcode::G_FMAXNUM:
1094 case TargetOpcode::G_FMAXIMUM:
1095 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1097 case TargetOpcode::G_FCOPYSIGN:
1098 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1100 case TargetOpcode::G_FCEIL:
1101 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1102 case TargetOpcode::G_FFLOOR:
1103 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1105 case TargetOpcode::G_FCOS:
1106 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1107 case TargetOpcode::G_FSIN:
1108 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1109 case TargetOpcode::G_FTAN:
1110 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1111 case TargetOpcode::G_FACOS:
1112 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1113 case TargetOpcode::G_FASIN:
1114 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1115 case TargetOpcode::G_FATAN:
1116 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1117 case TargetOpcode::G_FATAN2:
1118 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1119 case TargetOpcode::G_FCOSH:
1120 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1121 case TargetOpcode::G_FSINH:
1122 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1123 case TargetOpcode::G_FTANH:
1124 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1126 case TargetOpcode::G_STRICT_FSQRT:
1127 case TargetOpcode::G_FSQRT:
1128 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1130 case TargetOpcode::G_CTTZ:
1131 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1132 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1133 case TargetOpcode::G_CTLZ:
1134 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1135 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1137 case TargetOpcode::G_INTRINSIC_ROUND:
1138 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1139 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1140 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1141 case TargetOpcode::G_INTRINSIC_TRUNC:
1142 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1143 case TargetOpcode::G_FRINT:
1144 case TargetOpcode::G_FNEARBYINT:
1145 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1147 case TargetOpcode::G_SMULH:
1148 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1149 case TargetOpcode::G_UMULH:
1150 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1152 case TargetOpcode::G_SADDSAT:
1153 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1154 case TargetOpcode::G_UADDSAT:
1155 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1156 case TargetOpcode::G_SSUBSAT:
1157 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1158 case TargetOpcode::G_USUBSAT:
1159 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1161 case TargetOpcode::G_FFREXP:
1162 return selectFrexp(ResVReg, ResType,
I);
1164 case TargetOpcode::G_UADDO:
1165 return selectOverflowArith(ResVReg, ResType,
I,
1166 ResType->
getOpcode() == SPIRV::OpTypeVector
1167 ? SPIRV::OpIAddCarryV
1168 : SPIRV::OpIAddCarryS);
1169 case TargetOpcode::G_USUBO:
1170 return selectOverflowArith(ResVReg, ResType,
I,
1171 ResType->
getOpcode() == SPIRV::OpTypeVector
1172 ? SPIRV::OpISubBorrowV
1173 : SPIRV::OpISubBorrowS);
1174 case TargetOpcode::G_UMULO:
1175 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1176 case TargetOpcode::G_SMULO:
1177 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1179 case TargetOpcode::G_SEXT:
1180 return selectExt(ResVReg, ResType,
I,
true);
1181 case TargetOpcode::G_ANYEXT:
1182 case TargetOpcode::G_ZEXT:
1183 return selectExt(ResVReg, ResType,
I,
false);
1184 case TargetOpcode::G_TRUNC:
1185 return selectTrunc(ResVReg, ResType,
I);
1186 case TargetOpcode::G_FPTRUNC:
1187 case TargetOpcode::G_FPEXT:
1188 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1190 case TargetOpcode::G_PTRTOINT:
1191 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1192 case TargetOpcode::G_INTTOPTR:
1193 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1194 case TargetOpcode::G_BITCAST:
1195 return selectBitcast(ResVReg, ResType,
I);
1196 case TargetOpcode::G_ADDRSPACE_CAST:
1197 return selectAddrSpaceCast(ResVReg, ResType,
I);
1198 case TargetOpcode::G_PTR_ADD: {
1200 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1204 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1205 (*II).getOpcode() == TargetOpcode::COPY ||
1206 (*II).getOpcode() == SPIRV::OpVariable) &&
1207 getImm(
I.getOperand(2), MRI));
1209 bool IsGVInit =
false;
1213 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1214 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1215 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1216 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1226 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1239 "incompatible result and operand types in a bitcast");
1241 MachineInstrBuilder MIB =
1242 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1249 : SPIRV::OpInBoundsPtrAccessChain))
1253 .
addUse(
I.getOperand(2).getReg())
1256 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1260 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1262 .
addUse(
I.getOperand(2).getReg())
1271 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1274 .
addImm(
static_cast<uint32_t
>(
1275 SPIRV::Opcode::InBoundsPtrAccessChain))
1278 .
addUse(
I.getOperand(2).getReg());
1283 case TargetOpcode::G_ATOMICRMW_OR:
1284 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1285 case TargetOpcode::G_ATOMICRMW_ADD:
1286 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1287 case TargetOpcode::G_ATOMICRMW_AND:
1288 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1289 case TargetOpcode::G_ATOMICRMW_MAX:
1290 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1291 case TargetOpcode::G_ATOMICRMW_MIN:
1292 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1293 case TargetOpcode::G_ATOMICRMW_SUB:
1294 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1295 case TargetOpcode::G_ATOMICRMW_XOR:
1296 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1297 case TargetOpcode::G_ATOMICRMW_UMAX:
1298 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1299 case TargetOpcode::G_ATOMICRMW_UMIN:
1300 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1301 case TargetOpcode::G_ATOMICRMW_XCHG:
1302 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1303 case TargetOpcode::G_ATOMIC_CMPXCHG:
1304 return selectAtomicCmpXchg(ResVReg, ResType,
I);
1306 case TargetOpcode::G_ATOMICRMW_FADD:
1307 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1308 case TargetOpcode::G_ATOMICRMW_FSUB:
1310 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1311 ResType->
getOpcode() == SPIRV::OpTypeVector
1313 : SPIRV::OpFNegate);
1314 case TargetOpcode::G_ATOMICRMW_FMIN:
1315 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1316 case TargetOpcode::G_ATOMICRMW_FMAX:
1317 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1319 case TargetOpcode::G_FENCE:
1320 return selectFence(
I);
1322 case TargetOpcode::G_STACKSAVE:
1323 return selectStackSave(ResVReg, ResType,
I);
1324 case TargetOpcode::G_STACKRESTORE:
1325 return selectStackRestore(
I);
1327 case TargetOpcode::G_UNMERGE_VALUES:
1333 case TargetOpcode::G_TRAP:
1334 case TargetOpcode::G_UBSANTRAP:
1335 case TargetOpcode::DBG_LABEL:
1337 case TargetOpcode::G_DEBUGTRAP:
1338 return selectDebugTrap(ResVReg, ResType,
I);
1345bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1346 SPIRVTypeInst ResType,
1347 MachineInstr &
I)
const {
1348 unsigned Opcode = SPIRV::OpNop;
1355bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1356 SPIRVTypeInst ResType,
1358 GL::GLSLExtInst GLInst,
1359 bool setMIFlags,
bool useMISrc,
1362 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1363 std::string DiagMsg;
1364 raw_string_ostream OS(DiagMsg);
1365 I.print(OS,
true,
false,
false,
false);
1366 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1369 return selectExtInst(ResVReg, ResType,
I,
1370 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1371 setMIFlags, useMISrc, SrcRegs);
1374bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1375 SPIRVTypeInst ResType,
1377 CL::OpenCLExtInst CLInst,
1378 bool setMIFlags,
bool useMISrc,
1380 return selectExtInst(ResVReg, ResType,
I,
1381 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1382 setMIFlags, useMISrc, SrcRegs);
1385bool SPIRVInstructionSelector::selectExtInst(
1386 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1387 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1389 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1390 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1391 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1395bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1396 SPIRVTypeInst ResType,
1399 bool setMIFlags,
bool useMISrc,
1402 for (
const auto &[InstructionSet, Opcode] : Insts) {
1406 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1409 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1414 const unsigned NumOps =
I.getNumOperands();
1417 I.getOperand(Index).getType() ==
1418 MachineOperand::MachineOperandType::MO_IntrinsicID)
1421 MIB.
add(
I.getOperand(Index));
1433bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1434 SPIRVTypeInst ResType,
1435 MachineInstr &
I)
const {
1436 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1437 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1438 for (
const auto &Ex : ExtInsts) {
1439 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1440 uint32_t Opcode = Ex.second;
1444 MachineIRBuilder MIRBuilder(
I);
1447 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1452 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1455 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1458 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1461 .
addImm(
static_cast<uint32_t
>(Ex.first))
1463 .
add(
I.getOperand(2))
1467 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1468 .
addDef(
I.getOperand(1).getReg())
1477bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1478 SPIRVTypeInst ResType,
1479 MachineInstr &
I)
const {
1480 Register CosResVReg =
I.getOperand(1).getReg();
1481 unsigned SrcIdx =
I.getNumExplicitDefs();
1486 MachineIRBuilder MIRBuilder(
I);
1488 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1493 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1496 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1498 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1501 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1503 .
add(
I.getOperand(SrcIdx))
1506 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1514 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1517 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1519 .
add(
I.getOperand(SrcIdx))
1521 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1524 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1526 .
add(
I.getOperand(SrcIdx))
1533bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1534 SPIRVTypeInst ResType,
1536 std::vector<Register> Srcs,
1537 unsigned Opcode)
const {
1538 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1548bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1549 SPIRVTypeInst ResType,
1551 unsigned Opcode)
const {
1553 Register SrcReg =
I.getOperand(1).getReg();
1558 unsigned DefOpCode = DefIt->getOpcode();
1559 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1562 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1563 DefOpCode = VRD->getOpcode();
1565 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1566 DefOpCode == TargetOpcode::G_CONSTANT ||
1567 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1573 uint32_t SpecOpcode = 0;
1575 case SPIRV::OpConvertPtrToU:
1576 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1578 case SPIRV::OpConvertUToPtr:
1579 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1584 TII.get(SPIRV::OpSpecConstantOp))
1594 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1598bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1599 SPIRVTypeInst ResType,
1600 MachineInstr &
I)
const {
1601 Register OpReg =
I.getOperand(1).getReg();
1602 SPIRVTypeInst OpType =
1606 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1616 if (
MemOp->isVolatile())
1617 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1618 if (
MemOp->isNonTemporal())
1619 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1621 if (!ST->isShader() &&
MemOp->getAlign().value())
1622 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1626 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1627 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1631 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1633 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1637 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1641 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1643 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1655 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1657 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1659 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1663bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1664 SPIRVTypeInst ResType,
1665 MachineInstr &
I)
const {
1667 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1672 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1673 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1675 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1679 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1683 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1684 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1685 I.getDebugLoc(),
I);
1689 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1693 if (!
I.getNumMemOperands()) {
1694 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1696 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1699 MachineIRBuilder MIRBuilder(
I);
1706bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1708 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1709 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1714 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1715 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1720 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1724 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1725 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1726 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1727 TII.get(SPIRV::OpImageWrite))
1733 if (sampledTypeIsSignedInteger(LLVMHandleType))
1736 BMI.constrainAllUses(
TII,
TRI, RBI);
1742 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpStore))
1745 if (!
I.getNumMemOperands()) {
1746 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1748 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1751 MachineIRBuilder MIRBuilder(
I);
1758bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
1759 SPIRVTypeInst ResType,
1760 MachineInstr &
I)
const {
1761 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
1769 const Register PtrsReg =
I.getOperand(2).getReg();
1770 const uint32_t Alignment =
I.getOperand(3).getImm();
1771 const Register MaskReg =
I.getOperand(4).getReg();
1772 const Register PassthruReg =
I.getOperand(5).getReg();
1773 const Register AlignmentReg = buildI32Constant(Alignment,
I);
1777 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
1788bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
1789 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
1796 const Register ValuesReg =
I.getOperand(1).getReg();
1797 const Register PtrsReg =
I.getOperand(2).getReg();
1798 const uint32_t Alignment =
I.getOperand(3).getImm();
1799 const Register MaskReg =
I.getOperand(4).getReg();
1800 const Register AlignmentReg = buildI32Constant(Alignment,
I);
1804 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
1813bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
1814 const Twine &Msg)
const {
1815 const Function &
F =
I.getMF()->getFunction();
1816 F.getContext().diagnose(
1817 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
1821bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
1822 SPIRVTypeInst ResType,
1823 MachineInstr &
I)
const {
1824 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1826 "llvm.stacksave intrinsic: this instruction requires the following "
1827 "SPIR-V extension: SPV_INTEL_variable_length_array",
1830 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
1837bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
1838 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1840 "llvm.stackrestore intrinsic: this instruction requires the following "
1841 "SPIR-V extension: SPV_INTEL_variable_length_array",
1843 if (!
I.getOperand(0).isReg())
1846 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
1847 .
addUse(
I.getOperand(0).getReg())
1853SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
1854 MachineIRBuilder MIRBuilder(
I);
1855 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1862 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
1866 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1867 Type *ArrTy = ArrayType::get(ValTy, Num);
1869 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
1872 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
1879 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
1882 .
addImm(SPIRV::StorageClass::UniformConstant)
1893bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
1896 Register DstReg =
I.getOperand(0).getReg();
1901 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
1906 "Unable to determine pointee type size for OpCopyMemory");
1907 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
1908 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
1910 "OpCopyMemory requires the size to match the pointee type size");
1911 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
1914 if (
I.getNumMemOperands()) {
1915 MachineIRBuilder MIRBuilder(
I);
1922bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
1925 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
1926 .
addUse(
I.getOperand(0).getReg())
1928 .
addUse(
I.getOperand(2).getReg());
1929 if (
I.getNumMemOperands()) {
1930 MachineIRBuilder MIRBuilder(
I);
1937bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
1938 MachineInstr &
I)
const {
1939 Register SrcReg =
I.getOperand(1).getReg();
1940 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
1941 Register VarReg = getOrCreateMemSetGlobal(
I);
1944 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1946 ValTy,
I, SPIRV::StorageClass::UniformConstant);
1948 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
1952 if (!selectCopyMemory(
I, SrcReg))
1955 if (!selectCopyMemorySized(
I, SrcReg))
1958 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
1959 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
1964bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
1965 SPIRVTypeInst ResType,
1968 unsigned NegateOpcode)
const {
1970 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1973 Register ScopeReg = buildI32Constant(Scope,
I);
1975 Register Ptr =
I.getOperand(1).getReg();
1981 Register MemSemReg = buildI32Constant(MemSem ,
I);
1983 Register ValueReg =
I.getOperand(2).getReg();
1984 if (NegateOpcode != 0) {
1987 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
1992 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2003bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2004 unsigned ArgI =
I.getNumOperands() - 1;
2006 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2007 SPIRVTypeInst SrcType =
2009 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2011 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2013 SPIRVTypeInst ScalarType =
2016 unsigned CurrentIndex = 0;
2017 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2018 Register ResVReg =
I.getOperand(i).getReg();
2021 LLT ResLLT = MRI->
getType(ResVReg);
2027 ResType = ScalarType;
2033 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2036 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2042 for (
unsigned j = 0;
j < NumElements; ++
j) {
2043 MIB.
addImm(CurrentIndex + j);
2045 CurrentIndex += NumElements;
2049 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2061bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2064 Register MemSemReg = buildI32Constant(MemSem,
I);
2066 uint32_t
Scope =
static_cast<uint32_t
>(
2068 Register ScopeReg = buildI32Constant(Scope,
I);
2070 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2077bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2078 SPIRVTypeInst ResType,
2080 unsigned Opcode)
const {
2081 Type *ResTy =
nullptr;
2085 "Not enough info to select the arithmetic with overflow instruction");
2088 "with overflow instruction");
2094 MachineIRBuilder MIRBuilder(
I);
2096 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2097 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2103 Register ZeroReg = buildZerosVal(ResType,
I);
2108 if (ResName.
size() > 0)
2113 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2116 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2117 MIB.
addUse(
I.getOperand(i).getReg());
2122 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2123 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2125 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2126 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2133 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2134 .
addDef(
I.getOperand(1).getReg())
2142bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2143 SPIRVTypeInst ResType,
2144 MachineInstr &
I)
const {
2148 Register Ptr =
I.getOperand(2).getReg();
2151 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2154 ScopeReg = buildI32Constant(Scope,
I);
2156 unsigned ScSem =
static_cast<uint32_t
>(
2159 unsigned MemSemEq =
static_cast<uint32_t
>(
getMemSemantics(AO)) | ScSem;
2160 Register MemSemEqReg = buildI32Constant(MemSemEq,
I);
2162 unsigned MemSemNeq =
static_cast<uint32_t
>(
getMemSemantics(FO)) | ScSem;
2163 if (MemSemEq == MemSemNeq)
2164 MemSemNeqReg = MemSemEqReg;
2166 MemSemNeqReg = buildI32Constant(MemSemEq,
I);
2169 ScopeReg =
I.getOperand(5).getReg();
2170 MemSemEqReg =
I.getOperand(6).getReg();
2171 MemSemNeqReg =
I.getOperand(7).getReg();
2175 Register Val =
I.getOperand(4).getReg();
2179 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2198 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2205 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2217 case SPIRV::StorageClass::DeviceOnlyINTEL:
2218 case SPIRV::StorageClass::HostOnlyINTEL:
2227 bool IsGRef =
false;
2228 bool IsAllowedRefs =
2230 unsigned Opcode = It.getOpcode();
2231 if (Opcode == SPIRV::OpConstantComposite ||
2232 Opcode == SPIRV::OpVariable ||
2233 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2234 return IsGRef = true;
2235 return Opcode == SPIRV::OpName;
2237 return IsAllowedRefs && IsGRef;
2240Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2241 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2243 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2247SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2249 uint32_t Opcode)
const {
2250 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2251 TII.get(SPIRV::OpSpecConstantOp))
2259SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2260 SPIRVTypeInst SrcPtrTy)
const {
2261 SPIRVTypeInst GenericPtrTy =
2265 SPIRV::StorageClass::Generic),
2267 MachineFunction *MF =
I.getParent()->getParent();
2269 MachineInstrBuilder MIB = buildSpecConstantOp(
2271 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2281bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2282 SPIRVTypeInst ResType,
2283 MachineInstr &
I)
const {
2287 Register SrcPtr =
I.getOperand(1).getReg();
2291 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2292 ResType->
getOpcode() != SPIRV::OpTypePointer)
2293 return BuildCOPY(ResVReg, SrcPtr,
I);
2303 unsigned SpecOpcode =
2305 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2308 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2315 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2317 .constrainAllUses(
TII,
TRI, RBI);
2319 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2321 buildSpecConstantOp(
2323 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2324 .constrainAllUses(
TII,
TRI, RBI);
2331 return BuildCOPY(ResVReg, SrcPtr,
I);
2333 if ((SrcSC == SPIRV::StorageClass::Function &&
2334 DstSC == SPIRV::StorageClass::Private) ||
2335 (DstSC == SPIRV::StorageClass::Function &&
2336 SrcSC == SPIRV::StorageClass::Private))
2337 return BuildCOPY(ResVReg, SrcPtr,
I);
2341 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2344 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2347 SPIRVTypeInst GenericPtrTy =
2366 return selectUnOp(ResVReg, ResType,
I,
2367 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2369 return selectUnOp(ResVReg, ResType,
I,
2370 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2372 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2374 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2384 return SPIRV::OpFOrdEqual;
2386 return SPIRV::OpFOrdGreaterThanEqual;
2388 return SPIRV::OpFOrdGreaterThan;
2390 return SPIRV::OpFOrdLessThanEqual;
2392 return SPIRV::OpFOrdLessThan;
2394 return SPIRV::OpFOrdNotEqual;
2396 return SPIRV::OpOrdered;
2398 return SPIRV::OpFUnordEqual;
2400 return SPIRV::OpFUnordGreaterThanEqual;
2402 return SPIRV::OpFUnordGreaterThan;
2404 return SPIRV::OpFUnordLessThanEqual;
2406 return SPIRV::OpFUnordLessThan;
2408 return SPIRV::OpFUnordNotEqual;
2410 return SPIRV::OpUnordered;
2420 return SPIRV::OpIEqual;
2422 return SPIRV::OpINotEqual;
2424 return SPIRV::OpSGreaterThanEqual;
2426 return SPIRV::OpSGreaterThan;
2428 return SPIRV::OpSLessThanEqual;
2430 return SPIRV::OpSLessThan;
2432 return SPIRV::OpUGreaterThanEqual;
2434 return SPIRV::OpUGreaterThan;
2436 return SPIRV::OpULessThanEqual;
2438 return SPIRV::OpULessThan;
2447 return SPIRV::OpPtrEqual;
2449 return SPIRV::OpPtrNotEqual;
2460 return SPIRV::OpLogicalEqual;
2462 return SPIRV::OpLogicalNotEqual;
2496bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2497 SPIRVTypeInst ResType,
2499 unsigned OpAnyOrAll)
const {
2500 assert(
I.getNumOperands() == 3);
2501 assert(
I.getOperand(2).isReg());
2503 Register InputRegister =
I.getOperand(2).getReg();
2506 assert(InputType &&
"VReg has no type assigned");
2509 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2510 if (IsBoolTy && !IsVectorTy) {
2511 assert(ResVReg ==
I.getOperand(0).getReg());
2512 return BuildCOPY(ResVReg, InputRegister,
I);
2516 unsigned SpirvNotEqualId =
2517 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2519 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2524 IsBoolTy ? InputRegister
2532 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2534 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2551bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2552 SPIRVTypeInst ResType,
2553 MachineInstr &
I)
const {
2554 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2557bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2558 SPIRVTypeInst ResType,
2559 MachineInstr &
I)
const {
2560 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2564bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2565 SPIRVTypeInst ResType,
2566 MachineInstr &
I)
const {
2567 assert(
I.getNumOperands() == 4);
2568 assert(
I.getOperand(2).isReg());
2569 assert(
I.getOperand(3).isReg());
2571 [[maybe_unused]] SPIRVTypeInst VecType =
2576 "dot product requires a vector of at least 2 components");
2578 [[maybe_unused]] SPIRVTypeInst EltType =
2587 .
addUse(
I.getOperand(2).getReg())
2588 .
addUse(
I.getOperand(3).getReg())
2593bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2594 SPIRVTypeInst ResType,
2597 assert(
I.getNumOperands() == 4);
2598 assert(
I.getOperand(2).isReg());
2599 assert(
I.getOperand(3).isReg());
2602 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2606 .
addUse(
I.getOperand(2).getReg())
2607 .
addUse(
I.getOperand(3).getReg())
2614bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2615 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2616 assert(
I.getNumOperands() == 4);
2617 assert(
I.getOperand(2).isReg());
2618 assert(
I.getOperand(3).isReg());
2622 Register Vec0 =
I.getOperand(2).getReg();
2623 Register Vec1 =
I.getOperand(3).getReg();
2627 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2636 "dot product requires a vector of at least 2 components");
2639 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2649 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2660 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2672bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2673 SPIRVTypeInst ResType,
2674 MachineInstr &
I)
const {
2676 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2679 .
addUse(
I.getOperand(2).getReg())
2684bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2685 SPIRVTypeInst ResType,
2686 MachineInstr &
I)
const {
2688 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2691 .
addUse(
I.getOperand(2).getReg())
2696template <
bool Signed>
2697bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2698 SPIRVTypeInst ResType,
2699 MachineInstr &
I)
const {
2700 assert(
I.getNumOperands() == 5);
2701 assert(
I.getOperand(2).isReg());
2702 assert(
I.getOperand(3).isReg());
2703 assert(
I.getOperand(4).isReg());
2706 Register Acc =
I.getOperand(2).getReg();
2710 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2712 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
2717 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
2720 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2732template <
bool Signed>
2733bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
2734 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2735 assert(
I.getNumOperands() == 5);
2736 assert(
I.getOperand(2).isReg());
2737 assert(
I.getOperand(3).isReg());
2738 assert(
I.getOperand(4).isReg());
2741 Register Acc =
I.getOperand(2).getReg();
2747 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
2751 for (
unsigned i = 0; i < 4; i++) {
2774 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
2794 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2809bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
2810 SPIRVTypeInst ResType,
2811 MachineInstr &
I)
const {
2812 assert(
I.getNumOperands() == 3);
2813 assert(
I.getOperand(2).isReg());
2815 Register VZero = buildZerosValF(ResType,
I);
2816 Register VOne = buildOnesValF(ResType,
I);
2818 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
2821 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2823 .
addUse(
I.getOperand(2).getReg())
2830bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
2831 SPIRVTypeInst ResType,
2832 MachineInstr &
I)
const {
2833 assert(
I.getNumOperands() == 3);
2834 assert(
I.getOperand(2).isReg());
2836 Register InputRegister =
I.getOperand(2).getReg();
2838 auto &
DL =
I.getDebugLoc();
2848 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
2850 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
2858 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2863 if (NeedsConversion) {
2864 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
2875bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
2876 SPIRVTypeInst ResType,
2878 unsigned Opcode)
const {
2882 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2888 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
2889 BMI.addUse(
I.getOperand(J).getReg());
2896bool SPIRVInstructionSelector::selectWaveActiveCountBits(
2897 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2902 if (!selectWaveOpInst(BallotReg, BallotType,
I,
2903 SPIRV::OpGroupNonUniformBallot))
2908 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2913 .
addImm(SPIRV::GroupOperation::Reduce)
2922 if (
Type->getOpcode() != SPIRV::OpTypeVector)
2926 return Type->getOperand(2).getImm();
2929bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
2930 SPIRVTypeInst ResType,
2931 MachineInstr &
I)
const {
2936 Register InputReg =
I.getOperand(2).getReg();
2941 bool IsVector = NumElems > 1;
2944 SPIRVTypeInst ElemInputType = InputType;
2945 SPIRVTypeInst ElemBoolType = ResType;
2958 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
2959 SPIRV::OpGroupNonUniformAllEqual);
2964 ElementResults.
reserve(NumElems);
2966 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
2979 ElemInput = Extracted;
2985 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
2996 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3007bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3008 SPIRVTypeInst ResType,
3009 MachineInstr &
I)
const {
3011 assert(
I.getNumOperands() == 3);
3013 auto Op =
I.getOperand(2);
3025 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3047 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3051 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3058bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3059 SPIRVTypeInst ResType,
3061 bool IsUnsigned)
const {
3062 return selectWaveReduce(
3063 ResVReg, ResType,
I, IsUnsigned,
3064 [&](
Register InputRegister,
bool IsUnsigned) {
3065 const bool IsFloatTy =
3067 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3068 : SPIRV::OpGroupNonUniformSMax;
3069 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3073bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3074 SPIRVTypeInst ResType,
3076 bool IsUnsigned)
const {
3077 return selectWaveReduce(
3078 ResVReg, ResType,
I, IsUnsigned,
3079 [&](
Register InputRegister,
bool IsUnsigned) {
3080 const bool IsFloatTy =
3082 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3083 : SPIRV::OpGroupNonUniformSMin;
3084 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3088bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3089 SPIRVTypeInst ResType,
3090 MachineInstr &
I)
const {
3091 return selectWaveReduce(ResVReg, ResType,
I,
false,
3092 [&](
Register InputRegister,
bool IsUnsigned) {
3094 InputRegister, SPIRV::OpTypeFloat);
3095 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3096 : SPIRV::OpGroupNonUniformIAdd;
3100bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3101 SPIRVTypeInst ResType,
3102 MachineInstr &
I)
const {
3103 return selectWaveReduce(ResVReg, ResType,
I,
false,
3104 [&](
Register InputRegister,
bool IsUnsigned) {
3106 InputRegister, SPIRV::OpTypeFloat);
3107 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3108 : SPIRV::OpGroupNonUniformIMul;
3112template <
typename PickOpcodeFn>
3113bool SPIRVInstructionSelector::selectWaveReduce(
3114 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3115 PickOpcodeFn &&PickOpcode)
const {
3116 assert(
I.getNumOperands() == 3);
3117 assert(
I.getOperand(2).isReg());
3119 Register InputRegister =
I.getOperand(2).getReg();
3126 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3132 .
addImm(SPIRV::GroupOperation::Reduce)
3133 .
addUse(
I.getOperand(2).getReg())
3138bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3139 SPIRVTypeInst ResType,
3141 unsigned Opcode)
const {
3142 return selectWaveReduce(
3143 ResVReg, ResType,
I,
false,
3144 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3147bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3148 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3149 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3150 [&](
Register InputRegister,
bool IsUnsigned) {
3152 InputRegister, SPIRV::OpTypeFloat);
3154 ? SPIRV::OpGroupNonUniformFAdd
3155 : SPIRV::OpGroupNonUniformIAdd;
3159bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3160 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3161 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3162 [&](
Register InputRegister,
bool IsUnsigned) {
3164 InputRegister, SPIRV::OpTypeFloat);
3166 ? SPIRV::OpGroupNonUniformFMul
3167 : SPIRV::OpGroupNonUniformIMul;
3171template <
typename PickOpcodeFn>
3172bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3173 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3174 PickOpcodeFn &&PickOpcode)
const {
3175 assert(
I.getNumOperands() == 3);
3176 assert(
I.getOperand(2).isReg());
3178 Register InputRegister =
I.getOperand(2).getReg();
3185 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3191 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3192 .
addUse(
I.getOperand(2).getReg())
3197bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3198 SPIRVTypeInst ResType,
3201 assert(
I.getNumOperands() == 3);
3202 assert(
I.getOperand(2).isReg());
3204 Register InputRegister =
I.getOperand(2).getReg();
3210 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3221bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3222 SPIRVTypeInst ResType,
3227 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3232 : SPIRV::OpUConvert;
3236 ShiftOp = SPIRV::OpShiftRightLogicalV;
3241 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3242 TII.get(SPIRV::OpConstantComposite))
3245 for (
unsigned It = 0; It <
N; ++It)
3249 ShiftConst = CompositeReg;
3254 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3259 if (!selectBitreverse32(BitrevReg, Int32Type,
I, ExtReg))
3264 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3269 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3272bool SPIRVInstructionSelector::selectBitreverse32(
Register ResVReg,
3273 SPIRVTypeInst ResType,
3277 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3285bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3286 SPIRVTypeInst ResType,
3287 MachineInstr &
I)
const {
3288 Register OpReg =
I.getOperand(1).getReg();
3292 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3294 return selectBitreverse32(ResVReg, ResType,
I, OpReg);
3298bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3299 SPIRVTypeInst ResType,
3300 MachineInstr &
I)
const {
3306 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3308 Register OpReg =
I.getOperand(1).getReg();
3309 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
3310 if (
Def->getOpcode() == TargetOpcode::COPY)
3313 switch (
Def->getOpcode()) {
3314 case SPIRV::ASSIGN_TYPE:
3315 if (MachineInstr *AssignToDef =
3317 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3318 Reg =
Def->getOperand(2).getReg();
3321 case SPIRV::OpUndef:
3322 Reg =
Def->getOperand(1).getReg();
3325 unsigned DestOpCode;
3327 DestOpCode = SPIRV::OpConstantNull;
3329 DestOpCode = TargetOpcode::COPY;
3332 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3333 .
addDef(
I.getOperand(0).getReg())
3341bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3342 SPIRVTypeInst ResType,
3343 MachineInstr &
I)
const {
3345 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3347 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3351 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3356 for (
unsigned i =
I.getNumExplicitDefs();
3357 i <
I.getNumExplicitOperands() && IsConst; ++i)
3361 if (!IsConst &&
N < 2)
3363 "There must be at least two constituent operands in a vector");
3366 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3367 TII.get(IsConst ? SPIRV::OpConstantComposite
3368 : SPIRV::OpCompositeConstruct))
3371 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3372 MIB.
addUse(
I.getOperand(i).getReg());
3377bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3378 SPIRVTypeInst ResType,
3379 MachineInstr &
I)
const {
3381 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3383 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3389 if (!
I.getOperand(
OpIdx).isReg())
3396 if (!IsConst &&
N < 2)
3398 "There must be at least two constituent operands in a vector");
3401 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3402 TII.get(IsConst ? SPIRV::OpConstantComposite
3403 : SPIRV::OpCompositeConstruct))
3406 for (
unsigned i = 0; i <
N; ++i)
3412bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3413 SPIRVTypeInst ResType,
3414 MachineInstr &
I)
const {
3419 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3421 Opcode = SPIRV::OpDemoteToHelperInvocation;
3423 Opcode = SPIRV::OpKill;
3425 if (MachineInstr *NextI =
I.getNextNode()) {
3427 NextI->eraseFromParent();
3437bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3438 SPIRVTypeInst ResType,
unsigned CmpOpc,
3439 MachineInstr &
I)
const {
3440 Register Cmp0 =
I.getOperand(2).getReg();
3441 Register Cmp1 =
I.getOperand(3).getReg();
3444 "CMP operands should have the same type");
3445 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3455bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3456 SPIRVTypeInst ResType,
3457 MachineInstr &
I)
const {
3458 auto Pred =
I.getOperand(1).getPredicate();
3461 Register CmpOperand =
I.getOperand(2).getReg();
3468 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3472SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3473 SPIRVTypeInst ResType)
const {
3475 SPIRVTypeInst SpvI32Ty =
3478 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3485 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3488 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3491 .
addImm(APInt(32, Val).getZExtValue());
3493 GR.
add(ConstInt,
MI);
3498bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
3499 SPIRVTypeInst ResType,
3500 MachineInstr &
I)
const {
3502 return selectCmp(ResVReg, ResType, CmpOp,
I);
3505bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
3506 SPIRVTypeInst ResType,
3507 MachineInstr &
I)
const {
3509 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
3516 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
3517 ResType->
getOpcode() != SPIRV::OpTypeFloat)
3520 MachineIRBuilder MIRBuilder(
I);
3522 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
3528 "only float operands supported by GLSL extended math");
3531 MIRBuilder, SpirvScalarType);
3533 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
3534 ? SPIRV::OpVectorTimesScalar
3537 if (!selectOpWithSrcs(ArgReg, ResType,
I,
3538 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
3540 if (!selectExtInst(ResVReg, ResType,
I,
3541 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
3551Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
3552 MachineInstr &
I)
const {
3555 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3560bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
3566 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
3574 if (
Def->getOpcode() == SPIRV::OpConstantNull)
3577 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
3578 Def->getOpcode() == SPIRV::OpConstantI)
3591 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
3592 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
3594 Intrinsic::spv_const_composite)) {
3595 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
3596 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
3597 if (!IsZero(
Def->getOperand(i).getReg()))
3606Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
3607 MachineInstr &
I)
const {
3611 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3616Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
3617 MachineInstr &
I)
const {
3621 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3627 SPIRVTypeInst ResType,
3628 MachineInstr &
I)
const {
3632 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3637bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
3638 SPIRVTypeInst ResType,
3639 MachineInstr &
I)
const {
3640 Register SelectFirstArg =
I.getOperand(2).getReg();
3641 Register SelectSecondArg =
I.getOperand(3).getReg();
3650 SPIRV::OpTypeVector;
3657 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
3658 }
else if (IsPtrTy) {
3659 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
3661 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
3665 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
3666 }
else if (IsPtrTy) {
3667 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
3669 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3672 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3675 .
addUse(
I.getOperand(1).getReg())
3684bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
3685 SPIRVTypeInst ResType,
3687 MachineInstr &InsertAt,
3688 bool IsSigned)
const {
3690 Register ZeroReg = buildZerosVal(ResType, InsertAt);
3691 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
3692 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
3694 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3706bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
3707 SPIRVTypeInst ResType,
3708 MachineInstr &
I,
bool IsSigned,
3709 unsigned Opcode)
const {
3710 Register SrcReg =
I.getOperand(1).getReg();
3716 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
3721 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
3723 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
3726bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
3727 SPIRVTypeInst ResType, MachineInstr &
I,
3728 bool IsSigned)
const {
3729 Register SrcReg =
I.getOperand(1).getReg();
3731 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
3735 if (ResType == SrcType)
3736 return BuildCOPY(ResVReg, SrcReg,
I);
3738 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3739 return selectUnOp(ResVReg, ResType,
I, Opcode);
3742bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
3743 SPIRVTypeInst ResType,
3745 bool IsSigned)
const {
3746 MachineIRBuilder MIRBuilder(
I);
3747 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
3762 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
3765 .
addUse(
I.getOperand(1).getReg())
3766 .
addUse(
I.getOperand(2).getReg())
3772 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
3775 .
addUse(
I.getOperand(1).getReg())
3776 .
addUse(
I.getOperand(2).getReg())
3784 unsigned SelectOpcode =
3785 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
3790 .
addUse(buildOnesVal(
true, ResType,
I))
3791 .
addUse(buildZerosVal(ResType,
I))
3798 .
addUse(buildOnesVal(
false, ResType,
I))
3803bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
3806 SPIRVTypeInst IntTy,
3807 SPIRVTypeInst BoolTy)
const {
3810 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
3811 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
3813 Register One = buildOnesVal(
false, IntTy,
I);
3821 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
3830bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
3831 SPIRVTypeInst ResType,
3832 MachineInstr &
I)
const {
3833 Register IntReg =
I.getOperand(1).getReg();
3836 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
3837 if (ArgType == ResType)
3838 return BuildCOPY(ResVReg, IntReg,
I);
3840 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3841 return selectUnOp(ResVReg, ResType,
I, Opcode);
3844bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
3845 SPIRVTypeInst ResType,
3846 MachineInstr &
I)
const {
3847 unsigned Opcode =
I.getOpcode();
3848 unsigned TpOpcode = ResType->
getOpcode();
3850 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
3851 assert(Opcode == TargetOpcode::G_CONSTANT &&
3852 I.getOperand(1).getCImm()->isZero());
3853 MachineBasicBlock &DepMBB =
I.getMF()->front();
3856 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
3863 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
3866bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
3867 SPIRVTypeInst ResType,
3868 MachineInstr &
I)
const {
3869 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3876bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
3877 SPIRVTypeInst ResType,
3878 MachineInstr &
I)
const {
3880 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
3884 .
addUse(
I.getOperand(3).getReg())
3886 .
addUse(
I.getOperand(2).getReg());
3887 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
3893bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
3894 SPIRVTypeInst ResType,
3895 MachineInstr &
I)
const {
3896 Type *MaybeResTy =
nullptr;
3901 "Expected aggregate type for extractv instruction");
3903 SPIRV::AccessQualifier::ReadWrite,
false);
3907 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3910 .
addUse(
I.getOperand(2).getReg());
3911 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
3917bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
3918 SPIRVTypeInst ResType,
3919 MachineInstr &
I)
const {
3920 if (
getImm(
I.getOperand(4), MRI))
3921 return selectInsertVal(ResVReg, ResType,
I);
3923 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
3926 .
addUse(
I.getOperand(2).getReg())
3927 .
addUse(
I.getOperand(3).getReg())
3928 .
addUse(
I.getOperand(4).getReg())
3933bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
3934 SPIRVTypeInst ResType,
3935 MachineInstr &
I)
const {
3936 if (
getImm(
I.getOperand(3), MRI))
3937 return selectExtractVal(ResVReg, ResType,
I);
3939 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
3942 .
addUse(
I.getOperand(2).getReg())
3943 .
addUse(
I.getOperand(3).getReg())
3948bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
3949 SPIRVTypeInst ResType,
3950 MachineInstr &
I)
const {
3951 const bool IsGEPInBounds =
I.getOperand(2).getImm();
3957 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
3958 : SPIRV::OpAccessChain)
3959 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
3960 :
SPIRV::OpPtrAccessChain);
3962 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3966 .
addUse(
I.getOperand(3).getReg());
3968 (Opcode == SPIRV::OpPtrAccessChain ||
3969 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
3970 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
3971 "Cannot translate GEP to OpAccessChain. First index must be 0.");
3974 const unsigned StartingIndex =
3975 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
3978 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
3979 Res.addUse(
I.getOperand(i).getReg());
3980 Res.constrainAllUses(
TII,
TRI, RBI);
3985bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
3987 unsigned Lim =
I.getNumExplicitOperands();
3988 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
3989 Register OpReg =
I.getOperand(i).getReg();
3990 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
3992 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
3993 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
3994 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4001 MachineFunction *MF =
I.getMF();
4013 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4014 TII.get(SPIRV::OpSpecConstantOp))
4017 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4019 GR.
add(OpDefine, MIB);
4025bool SPIRVInstructionSelector::selectDerivativeInst(
4026 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4027 const unsigned DPdOpCode)
const {
4030 errorIfInstrOutsideShader(
I);
4035 Register SrcReg =
I.getOperand(2).getReg();
4040 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4043 .
addUse(
I.getOperand(2).getReg());
4045 MachineIRBuilder MIRBuilder(
I);
4048 if (componentCount != 1)
4052 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4056 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4061 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4066 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4074bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4075 SPIRVTypeInst ResType,
4076 MachineInstr &
I)
const {
4080 case Intrinsic::spv_load:
4081 return selectLoad(ResVReg, ResType,
I);
4082 case Intrinsic::spv_store:
4083 return selectStore(
I);
4084 case Intrinsic::spv_extractv:
4085 return selectExtractVal(ResVReg, ResType,
I);
4086 case Intrinsic::spv_insertv:
4087 return selectInsertVal(ResVReg, ResType,
I);
4088 case Intrinsic::spv_extractelt:
4089 return selectExtractElt(ResVReg, ResType,
I);
4090 case Intrinsic::spv_insertelt:
4091 return selectInsertElt(ResVReg, ResType,
I);
4092 case Intrinsic::spv_gep:
4093 return selectGEP(ResVReg, ResType,
I);
4094 case Intrinsic::spv_bitcast: {
4095 Register OpReg =
I.getOperand(2).getReg();
4096 SPIRVTypeInst OpType =
4100 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4102 case Intrinsic::spv_unref_global:
4103 case Intrinsic::spv_init_global: {
4104 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4109 Register GVarVReg =
MI->getOperand(0).getReg();
4110 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4115 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4117 MI->eraseFromParent();
4121 case Intrinsic::spv_undef: {
4122 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4128 case Intrinsic::spv_const_composite: {
4130 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4136 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4138 MachineIRBuilder MIR(
I);
4140 MIR, SPIRV::OpConstantComposite, 3,
4141 SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,
4143 for (
auto *Instr : Instructions) {
4144 Instr->setDebugLoc(
I.getDebugLoc());
4149 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4156 case Intrinsic::spv_assign_name: {
4157 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4158 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4159 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4160 i <
I.getNumExplicitOperands(); ++i) {
4161 MIB.
addImm(
I.getOperand(i).getImm());
4166 case Intrinsic::spv_switch: {
4167 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4168 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4169 if (
I.getOperand(i).isReg())
4170 MIB.
addReg(
I.getOperand(i).getReg());
4171 else if (
I.getOperand(i).isCImm())
4172 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4173 else if (
I.getOperand(i).isMBB())
4174 MIB.
addMBB(
I.getOperand(i).getMBB());
4181 case Intrinsic::spv_loop_merge: {
4182 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4183 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4184 if (
I.getOperand(i).isMBB())
4185 MIB.
addMBB(
I.getOperand(i).getMBB());
4192 case Intrinsic::spv_loop_control_intel: {
4194 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4195 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4200 case Intrinsic::spv_selection_merge: {
4202 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4203 assert(
I.getOperand(1).isMBB() &&
4204 "operand 1 to spv_selection_merge must be a basic block");
4205 MIB.
addMBB(
I.getOperand(1).getMBB());
4206 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4210 case Intrinsic::spv_cmpxchg:
4211 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4212 case Intrinsic::spv_unreachable:
4213 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4216 case Intrinsic::spv_alloca:
4217 return selectFrameIndex(ResVReg, ResType,
I);
4218 case Intrinsic::spv_alloca_array:
4219 return selectAllocaArray(ResVReg, ResType,
I);
4220 case Intrinsic::spv_assume:
4222 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4223 .
addUse(
I.getOperand(1).getReg())
4228 case Intrinsic::spv_expect:
4230 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4233 .
addUse(
I.getOperand(2).getReg())
4234 .
addUse(
I.getOperand(3).getReg())
4239 case Intrinsic::arithmetic_fence:
4240 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4241 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4244 .
addUse(
I.getOperand(2).getReg())
4248 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4250 case Intrinsic::spv_thread_id:
4256 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4258 case Intrinsic::spv_thread_id_in_group:
4264 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4266 case Intrinsic::spv_group_id:
4272 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4274 case Intrinsic::spv_flattened_thread_id_in_group:
4281 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4283 case Intrinsic::spv_workgroup_size:
4284 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4286 case Intrinsic::spv_global_size:
4287 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4289 case Intrinsic::spv_global_offset:
4290 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4292 case Intrinsic::spv_num_workgroups:
4293 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4295 case Intrinsic::spv_subgroup_size:
4296 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4298 case Intrinsic::spv_num_subgroups:
4299 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4301 case Intrinsic::spv_subgroup_id:
4302 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4303 case Intrinsic::spv_subgroup_local_invocation_id:
4304 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4305 ResVReg, ResType,
I);
4306 case Intrinsic::spv_subgroup_max_size:
4307 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4309 case Intrinsic::spv_fdot:
4310 return selectFloatDot(ResVReg, ResType,
I);
4311 case Intrinsic::spv_udot:
4312 case Intrinsic::spv_sdot:
4313 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4315 return selectIntegerDot(ResVReg, ResType,
I,
4316 IID == Intrinsic::spv_sdot);
4317 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4318 case Intrinsic::spv_dot4add_i8packed:
4319 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4321 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4322 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4323 case Intrinsic::spv_dot4add_u8packed:
4324 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4326 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4327 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4328 case Intrinsic::spv_all:
4329 return selectAll(ResVReg, ResType,
I);
4330 case Intrinsic::spv_any:
4331 return selectAny(ResVReg, ResType,
I);
4332 case Intrinsic::spv_cross:
4333 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4334 case Intrinsic::spv_distance:
4335 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4336 case Intrinsic::spv_lerp:
4337 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4338 case Intrinsic::spv_length:
4339 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4340 case Intrinsic::spv_degrees:
4341 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4342 case Intrinsic::spv_faceforward:
4343 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4344 case Intrinsic::spv_frac:
4345 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4346 case Intrinsic::spv_isinf:
4347 return selectOpIsInf(ResVReg, ResType,
I);
4348 case Intrinsic::spv_isnan:
4349 return selectOpIsNan(ResVReg, ResType,
I);
4350 case Intrinsic::spv_normalize:
4351 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4352 case Intrinsic::spv_refract:
4353 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4354 case Intrinsic::spv_reflect:
4355 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4356 case Intrinsic::spv_rsqrt:
4357 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4358 case Intrinsic::spv_sign:
4359 return selectSign(ResVReg, ResType,
I);
4360 case Intrinsic::spv_smoothstep:
4361 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4362 case Intrinsic::spv_firstbituhigh:
4363 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4364 case Intrinsic::spv_firstbitshigh:
4365 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4366 case Intrinsic::spv_firstbitlow:
4367 return selectFirstBitLow(ResVReg, ResType,
I);
4368 case Intrinsic::spv_group_memory_barrier_with_group_sync: {
4370 buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent,
I);
4371 Register ScopeReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
4373 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpControlBarrier))
4380 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4381 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4382 SPIRV::StorageClass::StorageClass ResSC =
4386 "Generic storage class");
4387 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4395 case Intrinsic::spv_lifetime_start:
4396 case Intrinsic::spv_lifetime_end: {
4397 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
4398 : SPIRV::OpLifetimeStop;
4399 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
4400 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
4409 case Intrinsic::spv_saturate:
4410 return selectSaturate(ResVReg, ResType,
I);
4411 case Intrinsic::spv_nclamp:
4412 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
4413 case Intrinsic::spv_uclamp:
4414 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
4415 case Intrinsic::spv_sclamp:
4416 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
4417 case Intrinsic::spv_subgroup_prefix_bit_count:
4418 return selectWavePrefixBitCount(ResVReg, ResType,
I);
4419 case Intrinsic::spv_wave_active_countbits:
4420 return selectWaveActiveCountBits(ResVReg, ResType,
I);
4421 case Intrinsic::spv_wave_all_equal:
4422 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
4423 case Intrinsic::spv_wave_all:
4424 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
4425 case Intrinsic::spv_wave_any:
4426 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
4427 case Intrinsic::spv_subgroup_ballot:
4428 return selectWaveOpInst(ResVReg, ResType,
I,
4429 SPIRV::OpGroupNonUniformBallot);
4430 case Intrinsic::spv_wave_is_first_lane:
4431 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
4432 case Intrinsic::spv_wave_reduce_or:
4433 return selectWaveReduceOp(ResVReg, ResType,
I,
4434 SPIRV::OpGroupNonUniformBitwiseOr);
4435 case Intrinsic::spv_wave_reduce_xor:
4436 return selectWaveReduceOp(ResVReg, ResType,
I,
4437 SPIRV::OpGroupNonUniformBitwiseXor);
4438 case Intrinsic::spv_wave_reduce_umax:
4439 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
4440 case Intrinsic::spv_wave_reduce_max:
4441 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
4442 case Intrinsic::spv_wave_reduce_umin:
4443 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
4444 case Intrinsic::spv_wave_reduce_min:
4445 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
4446 case Intrinsic::spv_wave_reduce_sum:
4447 return selectWaveReduceSum(ResVReg, ResType,
I);
4448 case Intrinsic::spv_wave_product:
4449 return selectWaveReduceProduct(ResVReg, ResType,
I);
4450 case Intrinsic::spv_wave_readlane:
4451 return selectWaveOpInst(ResVReg, ResType,
I,
4452 SPIRV::OpGroupNonUniformShuffle);
4453 case Intrinsic::spv_wave_prefix_sum:
4454 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
4455 case Intrinsic::spv_wave_prefix_product:
4456 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
4457 case Intrinsic::spv_quad_read_across_x: {
4458 return selectQuadSwap(ResVReg, ResType,
I, 0);
4460 case Intrinsic::spv_step:
4461 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
4462 case Intrinsic::spv_radians:
4463 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
4467 case Intrinsic::instrprof_increment:
4468 case Intrinsic::instrprof_increment_step:
4469 case Intrinsic::instrprof_value_profile:
4472 case Intrinsic::spv_value_md:
4474 case Intrinsic::spv_resource_handlefrombinding: {
4475 return selectHandleFromBinding(ResVReg, ResType,
I);
4477 case Intrinsic::spv_resource_counterhandlefrombinding:
4478 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
4479 case Intrinsic::spv_resource_updatecounter:
4480 return selectUpdateCounter(ResVReg, ResType,
I);
4481 case Intrinsic::spv_resource_store_typedbuffer: {
4482 return selectImageWriteIntrinsic(
I);
4484 case Intrinsic::spv_resource_load_typedbuffer: {
4485 return selectReadImageIntrinsic(ResVReg, ResType,
I);
4487 case Intrinsic::spv_resource_load_level: {
4488 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
4490 case Intrinsic::spv_resource_sample:
4491 case Intrinsic::spv_resource_sample_clamp:
4492 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
4493 case Intrinsic::spv_resource_samplebias:
4494 case Intrinsic::spv_resource_samplebias_clamp:
4495 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
4496 case Intrinsic::spv_resource_samplegrad:
4497 case Intrinsic::spv_resource_samplegrad_clamp:
4498 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
4499 case Intrinsic::spv_resource_samplelevel:
4500 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
4501 case Intrinsic::spv_resource_samplecmp:
4502 case Intrinsic::spv_resource_samplecmp_clamp:
4503 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
4504 case Intrinsic::spv_resource_samplecmplevelzero:
4505 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
4506 case Intrinsic::spv_resource_gather:
4507 case Intrinsic::spv_resource_gather_cmp:
4508 return selectGatherIntrinsic(ResVReg, ResType,
I);
4509 case Intrinsic::spv_resource_getpointer: {
4510 return selectResourceGetPointer(ResVReg, ResType,
I);
4512 case Intrinsic::spv_pushconstant_getpointer: {
4513 return selectPushConstantGetPointer(ResVReg, ResType,
I);
4515 case Intrinsic::spv_discard: {
4516 return selectDiscard(ResVReg, ResType,
I);
4518 case Intrinsic::spv_resource_nonuniformindex: {
4519 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
4521 case Intrinsic::spv_unpackhalf2x16: {
4522 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
4524 case Intrinsic::spv_packhalf2x16: {
4525 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
4527 case Intrinsic::spv_ddx:
4528 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
4529 case Intrinsic::spv_ddy:
4530 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
4531 case Intrinsic::spv_ddx_coarse:
4532 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
4533 case Intrinsic::spv_ddy_coarse:
4534 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
4535 case Intrinsic::spv_ddx_fine:
4536 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
4537 case Intrinsic::spv_ddy_fine:
4538 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
4539 case Intrinsic::spv_fwidth:
4540 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
4541 case Intrinsic::spv_masked_gather:
4542 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4543 return selectMaskedGather(ResVReg, ResType,
I);
4544 return diagnoseUnsupported(
4545 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
4546 case Intrinsic::spv_masked_scatter:
4547 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4548 return selectMaskedScatter(
I);
4549 return diagnoseUnsupported(
4550 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
4552 std::string DiagMsg;
4553 raw_string_ostream OS(DiagMsg);
4555 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
4562bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
4563 SPIRVTypeInst ResType,
4564 MachineInstr &
I)
const {
4567 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
4574bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
4575 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4577 assert(Intr.getIntrinsicID() ==
4578 Intrinsic::spv_resource_counterhandlefrombinding);
4581 Register MainHandleReg = Intr.getOperand(2).getReg();
4583 assert(MainHandleDef->getIntrinsicID() ==
4584 Intrinsic::spv_resource_handlefrombinding);
4588 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
4589 Register IndexReg = MainHandleDef->getOperand(5).getReg();
4590 std::string CounterName =
4595 MachineIRBuilder MIRBuilder(
I);
4597 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
4599 ArraySize, IndexReg, CounterName, MIRBuilder);
4601 return BuildCOPY(ResVReg, CounterVarReg,
I);
4604bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
4605 SPIRVTypeInst ResType,
4606 MachineInstr &
I)
const {
4608 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
4610 Register CounterHandleReg = Intr.getOperand(2).getReg();
4611 Register IncrReg = Intr.getOperand(3).getReg();
4618 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
4619 assert(CounterVarPointeeType &&
4620 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
4621 "Counter variable must be a struct");
4623 SPIRV::StorageClass::StorageBuffer &&
4624 "Counter variable must be in the storage buffer storage class");
4626 "Counter variable must have exactly 1 member in the struct");
4627 const SPIRVTypeInst MemberType =
4630 "Counter variable struct must have a single i32 member");
4634 MachineIRBuilder MIRBuilder(
I);
4636 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
4639 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
4645 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4648 .
addUse(CounterHandleReg)
4655 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
4658 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
4661 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
4670 return BuildCOPY(ResVReg, AtomicRes,
I);
4678 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
4686bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
4687 SPIRVTypeInst ResType,
4688 MachineInstr &
I)
const {
4696 Register ImageReg =
I.getOperand(2).getReg();
4704 Register IdxReg =
I.getOperand(3).getReg();
4706 MachineInstr &Pos =
I;
4708 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
4712bool SPIRVInstructionSelector::generateSampleImage(
4715 DebugLoc Loc, MachineInstr &Pos)
const {
4726 if (!loadHandleBeforePosition(NewSamplerReg,
4732 MachineIRBuilder MIRBuilder(Pos);
4745 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
4746 ImOps.Lod.has_value();
4747 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
4748 : SPIRV::OpImageSampleImplicitLod;
4750 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
4751 : SPIRV::OpImageSampleDrefImplicitLod;
4760 MIB.
addUse(*ImOps.Compare);
4762 uint32_t ImageOperands = 0;
4764 ImageOperands |= SPIRV::ImageOperand::Bias;
4766 ImageOperands |= SPIRV::ImageOperand::Lod;
4767 if (ImOps.GradX && ImOps.GradY)
4768 ImageOperands |= SPIRV::ImageOperand::Grad;
4769 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
4771 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4774 "Non-constant offsets are not supported in sample instructions.");
4778 ImageOperands |= SPIRV::ImageOperand::MinLod;
4780 if (ImageOperands != 0) {
4781 MIB.
addImm(ImageOperands);
4782 if (ImageOperands & SPIRV::ImageOperand::Bias)
4784 if (ImageOperands & SPIRV::ImageOperand::Lod)
4786 if (ImageOperands & SPIRV::ImageOperand::Grad) {
4787 MIB.
addUse(*ImOps.GradX);
4788 MIB.
addUse(*ImOps.GradY);
4791 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4792 MIB.
addUse(*ImOps.Offset);
4793 if (ImageOperands & SPIRV::ImageOperand::MinLod)
4794 MIB.
addUse(*ImOps.MinLod);
4801bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
4802 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4803 Register ImageReg =
I.getOperand(2).getReg();
4804 Register SamplerReg =
I.getOperand(3).getReg();
4805 Register CoordinateReg =
I.getOperand(4).getReg();
4806 ImageOperands ImOps;
4807 if (
I.getNumOperands() > 5)
4808 ImOps.Offset =
I.getOperand(5).getReg();
4809 if (
I.getNumOperands() > 6)
4810 ImOps.MinLod =
I.getOperand(6).getReg();
4811 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4812 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4815bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
4816 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4817 Register ImageReg =
I.getOperand(2).getReg();
4818 Register SamplerReg =
I.getOperand(3).getReg();
4819 Register CoordinateReg =
I.getOperand(4).getReg();
4820 ImageOperands ImOps;
4821 ImOps.Bias =
I.getOperand(5).getReg();
4822 if (
I.getNumOperands() > 6)
4823 ImOps.Offset =
I.getOperand(6).getReg();
4824 if (
I.getNumOperands() > 7)
4825 ImOps.MinLod =
I.getOperand(7).getReg();
4826 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4827 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4830bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
4831 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4832 Register ImageReg =
I.getOperand(2).getReg();
4833 Register SamplerReg =
I.getOperand(3).getReg();
4834 Register CoordinateReg =
I.getOperand(4).getReg();
4835 ImageOperands ImOps;
4836 ImOps.GradX =
I.getOperand(5).getReg();
4837 ImOps.GradY =
I.getOperand(6).getReg();
4838 if (
I.getNumOperands() > 7)
4839 ImOps.Offset =
I.getOperand(7).getReg();
4840 if (
I.getNumOperands() > 8)
4841 ImOps.MinLod =
I.getOperand(8).getReg();
4842 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4843 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4846bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
4847 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4848 Register ImageReg =
I.getOperand(2).getReg();
4849 Register SamplerReg =
I.getOperand(3).getReg();
4850 Register CoordinateReg =
I.getOperand(4).getReg();
4851 ImageOperands ImOps;
4852 ImOps.Lod =
I.getOperand(5).getReg();
4853 if (
I.getNumOperands() > 6)
4854 ImOps.Offset =
I.getOperand(6).getReg();
4855 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4856 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4859bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
4860 SPIRVTypeInst ResType,
4861 MachineInstr &
I)
const {
4862 Register ImageReg =
I.getOperand(2).getReg();
4863 Register SamplerReg =
I.getOperand(3).getReg();
4864 Register CoordinateReg =
I.getOperand(4).getReg();
4865 ImageOperands ImOps;
4866 ImOps.Compare =
I.getOperand(5).getReg();
4867 if (
I.getNumOperands() > 6)
4868 ImOps.Offset =
I.getOperand(6).getReg();
4869 if (
I.getNumOperands() > 7)
4870 ImOps.MinLod =
I.getOperand(7).getReg();
4871 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4872 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4875bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
4876 SPIRVTypeInst ResType,
4877 MachineInstr &
I)
const {
4878 Register ImageReg =
I.getOperand(2).getReg();
4879 Register CoordinateReg =
I.getOperand(3).getReg();
4880 Register LodReg =
I.getOperand(4).getReg();
4882 ImageOperands ImOps;
4884 if (
I.getNumOperands() > 5)
4885 ImOps.Offset =
I.getOperand(5).getReg();
4897 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
4898 I.getDebugLoc(),
I, &ImOps);
4901bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
4902 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4903 Register ImageReg =
I.getOperand(2).getReg();
4904 Register SamplerReg =
I.getOperand(3).getReg();
4905 Register CoordinateReg =
I.getOperand(4).getReg();
4906 ImageOperands ImOps;
4907 ImOps.Compare =
I.getOperand(5).getReg();
4908 if (
I.getNumOperands() > 6)
4909 ImOps.Offset =
I.getOperand(6).getReg();
4912 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4913 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4916bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
4917 SPIRVTypeInst ResType,
4918 MachineInstr &
I)
const {
4919 Register ImageReg =
I.getOperand(2).getReg();
4920 Register SamplerReg =
I.getOperand(3).getReg();
4921 Register CoordinateReg =
I.getOperand(4).getReg();
4924 "ImageReg is not an image type.");
4929 ComponentOrCompareReg =
I.getOperand(5).getReg();
4930 OffsetReg =
I.getOperand(6).getReg();
4933 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
4937 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
4938 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
4939 Dim != SPIRV::Dim::DIM_Rect) {
4941 "Gather operations are only supported for 2D, Cube, and Rect images.");
4948 if (!loadHandleBeforePosition(
4953 MachineIRBuilder MIRBuilder(
I);
4954 SPIRVTypeInst SampledImageType =
4959 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
4967 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
4969 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
4971 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4976 .
addUse(ComponentOrCompareReg);
4978 uint32_t ImageOperands = 0;
4979 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
4980 if (Dim == SPIRV::Dim::DIM_Cube) {
4982 "Gather operations with offset are not supported for Cube images.");
4986 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4988 ImageOperands |= SPIRV::ImageOperand::Offset;
4992 if (ImageOperands != 0) {
4993 MIB.
addImm(ImageOperands);
4995 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5003bool SPIRVInstructionSelector::generateImageReadOrFetch(
5006 const ImageOperands *ImOps)
const {
5009 "ImageReg is not an image type.");
5011 bool IsSignedInteger =
5016 bool IsFetch = (SampledOp.getImm() == 1);
5018 auto AddOperands = [&](MachineInstrBuilder &MIB) {
5019 uint32_t ImageOperandsMask = 0;
5020 if (IsSignedInteger)
5021 ImageOperandsMask |= 0x1000;
5023 if (IsFetch && ImOps) {
5025 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
5026 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
5028 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
5030 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5034 if (ImageOperandsMask != 0) {
5035 MIB.
addImm(ImageOperandsMask);
5036 if (IsFetch && ImOps) {
5039 if (ImOps->Offset &&
5040 (ImageOperandsMask &
5041 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5042 MIB.
addUse(*ImOps->Offset);
5048 if (ResultSize == 4) {
5051 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5058 BMI.constrainAllUses(
TII,
TRI, RBI);
5062 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5066 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5072 BMI.constrainAllUses(
TII,
TRI, RBI);
5074 if (ResultSize == 1) {
5083 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
5086bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
5087 SPIRVTypeInst ResType,
5088 MachineInstr &
I)
const {
5089 Register ResourcePtr =
I.getOperand(2).getReg();
5091 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
5100 MachineIRBuilder MIRBuilder(
I);
5102 Register IndexReg =
I.getOperand(3).getReg();
5105 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5115bool SPIRVInstructionSelector::selectPushConstantGetPointer(
5116 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5121bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
5122 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5123 Register ObjReg =
I.getOperand(2).getReg();
5124 if (!BuildCOPY(ResVReg, ObjReg,
I))
5134 decorateUsesAsNonUniform(ResVReg);
5138void SPIRVInstructionSelector::decorateUsesAsNonUniform(
5141 while (WorkList.
size() > 0) {
5145 bool IsDecorated =
false;
5147 if (
Use.getOpcode() == SPIRV::OpDecorate &&
5148 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
5154 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
5156 if (ResultReg == CurrentReg)
5164 SPIRV::Decoration::NonUniformEXT, {});
5169bool SPIRVInstructionSelector::extractSubvector(
5171 MachineInstr &InsertionPoint)
const {
5173 [[maybe_unused]] uint64_t InputSize =
5176 assert(InputSize > 1 &&
"The input must be a vector.");
5177 assert(ResultSize > 1 &&
"The result must be a vector.");
5178 assert(ResultSize < InputSize &&
5179 "Cannot extract more element than there are in the input.");
5182 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
5183 for (uint64_t
I = 0;
I < ResultSize;
I++) {
5186 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5195 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
5197 TII.get(SPIRV::OpCompositeConstruct))
5201 for (
Register ComponentReg : ComponentRegisters)
5202 MIB.
addUse(ComponentReg);
5207bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
5208 MachineInstr &
I)
const {
5215 Register ImageReg =
I.getOperand(1).getReg();
5223 Register CoordinateReg =
I.getOperand(2).getReg();
5224 Register DataReg =
I.getOperand(3).getReg();
5227 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
5235Register SPIRVInstructionSelector::buildPointerToResource(
5236 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
5237 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
5238 StringRef Name, MachineIRBuilder MIRBuilder)
const {
5240 if (ArraySize == 1) {
5241 SPIRVTypeInst PtrType =
5244 "SpirvResType did not have an explicit layout.");
5249 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
5250 SPIRVTypeInst VarPointerType =
5253 VarPointerType, Set,
Binding, Name, MIRBuilder);
5255 SPIRVTypeInst ResPointerType =
5268bool SPIRVInstructionSelector::selectFirstBitSet16(
5269 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5270 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
5272 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
5276 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
5279bool SPIRVInstructionSelector::selectFirstBitSet32(
5281 unsigned BitSetOpcode)
const {
5282 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5285 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5292bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
5294 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5301 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
5303 MachineIRBuilder MIRBuilder(
I);
5306 SPIRVTypeInst I64x2Type =
5308 SPIRVTypeInst Vec2ResType =
5311 std::vector<Register> PartialRegs;
5314 unsigned CurrentComponent = 0;
5315 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
5321 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5322 TII.get(SPIRV::OpVectorShuffle))
5327 .
addImm(CurrentComponent)
5328 .
addImm(CurrentComponent + 1);
5335 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType,
I, BitSetResult,
5336 BitSetOpcode, SwapPrimarySide))
5339 PartialRegs.push_back(SubVecBitSetReg);
5343 if (CurrentComponent != ComponentCount) {
5349 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
5350 SPIRV::OpVectorExtractDynamic))
5356 if (!selectFirstBitSet64(FinalElemBitSetReg,
BaseType,
I, FinalElemReg,
5357 BitSetOpcode, SwapPrimarySide))
5360 PartialRegs.push_back(FinalElemBitSetReg);
5365 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
5366 SPIRV::OpCompositeConstruct);
5369bool SPIRVInstructionSelector::selectFirstBitSet64(
5371 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5384 if (ComponentCount > 2) {
5385 return selectFirstBitSet64Overflow(ResVReg, ResType,
I, SrcReg,
5386 BitSetOpcode, SwapPrimarySide);
5390 MachineIRBuilder MIRBuilder(
I);
5392 BaseType, 2 * ComponentCount, MIRBuilder,
false);
5396 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
5402 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
5409 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
5412 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntZero},
5413 SPIRV::OpVectorExtractDynamic))
5415 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntOne},
5416 SPIRV::OpVectorExtractDynamic))
5420 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5421 TII.get(SPIRV::OpVectorShuffle))
5429 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
5435 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5436 TII.get(SPIRV::OpVectorShuffle))
5444 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
5464 SelectOp = SPIRV::OpSelectSISCond;
5465 AddOp = SPIRV::OpIAddS;
5473 SelectOp = SPIRV::OpSelectVIVCond;
5474 AddOp = SPIRV::OpIAddV;
5484 if (SwapPrimarySide) {
5485 PrimaryReg = LowReg;
5486 SecondaryReg = HighReg;
5487 PrimaryShiftReg = Reg0;
5488 SecondaryShiftReg = Reg32;
5493 if (!selectOpWithSrcs(BReg, BoolType,
I, {PrimaryReg, NegOneReg},
5499 if (!selectOpWithSrcs(TmpReg, ResType,
I, {
BReg, SecondaryReg, PrimaryReg},
5505 if (!selectOpWithSrcs(ValReg, ResType,
I,
5506 {
BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))
5509 return selectOpWithSrcs(ResVReg, ResType,
I, {ValReg, TmpReg}, AddOp);
5512bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
5513 SPIRVTypeInst ResType,
5515 bool IsSigned)
const {
5517 Register OpReg =
I.getOperand(2).getReg();
5520 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5521 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
5525 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5527 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5529 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5533 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
5537bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
5538 SPIRVTypeInst ResType,
5539 MachineInstr &
I)
const {
5541 Register OpReg =
I.getOperand(2).getReg();
5546 unsigned ExtendOpcode = SPIRV::OpUConvert;
5547 unsigned BitSetOpcode = GL::FindILsb;
5551 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5553 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5555 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5562bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
5563 SPIRVTypeInst ResType,
5564 MachineInstr &
I)
const {
5568 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
5571 .
addUse(
I.getOperand(2).getReg())
5574 unsigned Alignment =
I.getOperand(3).getImm();
5580bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
5581 SPIRVTypeInst ResType,
5582 MachineInstr &
I)
const {
5586 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
5589 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
5592 unsigned Alignment =
I.getOperand(2).getImm();
5599bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
5604 const MachineInstr *PrevI =
I.getPrevNode();
5606 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
5610 .
addMBB(
I.getOperand(0).getMBB())
5615 .
addMBB(
I.getOperand(0).getMBB())
5620bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
5631 const MachineInstr *NextI =
I.getNextNode();
5633 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
5639 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
5641 .
addUse(
I.getOperand(0).getReg())
5642 .
addMBB(
I.getOperand(1).getMBB())
5648bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
5649 MachineInstr &
I)
const {
5651 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
5653 const unsigned NumOps =
I.getNumOperands();
5654 for (
unsigned i = 1; i <
NumOps; i += 2) {
5655 MIB.
addUse(
I.getOperand(i + 0).getReg());
5656 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
5662bool SPIRVInstructionSelector::selectGlobalValue(
5663 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
5665 MachineIRBuilder MIRBuilder(
I);
5666 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
5669 std::string GlobalIdent;
5671 unsigned &
ID = UnnamedGlobalIDs[GV];
5673 ID = UnnamedGlobalIDs.
size();
5674 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
5700 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
5707 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
5712 MachineInstrBuilder MIB1 =
5713 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5716 MachineInstrBuilder MIB2 =
5718 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
5722 GR.
add(ConstVal, MIB2);
5730 MachineInstrBuilder MIB3 =
5731 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5734 GR.
add(ConstVal, MIB3);
5738 assert(NewReg != ResVReg);
5739 return BuildCOPY(ResVReg, NewReg,
I);
5749 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
5755 SPIRVTypeInst ResType =
5759 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
5764 if (
GlobalVar->isExternallyInitialized() &&
5765 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
5766 constexpr unsigned ReadWriteINTEL = 3u;
5769 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
5775bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
5776 SPIRVTypeInst ResType,
5777 MachineInstr &
I)
const {
5779 return selectExtInst(ResVReg, ResType,
I, CL::log10);
5787 MachineIRBuilder MIRBuilder(
I);
5792 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5795 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5797 .
add(
I.getOperand(1))
5802 ResType->
getOpcode() == SPIRV::OpTypeFloat);
5804 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
5812 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
5813 ? SPIRV::OpVectorTimesScalar
5824bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
5825 SPIRVTypeInst ResType,
5826 MachineInstr &
I)
const {
5829 return selectExtInst(ResVReg, ResType,
I, CL::pown);
5835 Register ExpReg =
I.getOperand(2).getReg();
5837 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
5838 SPIRV::OpConvertSToF))
5840 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
5847bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
5848 SPIRVTypeInst ResType,
5849 MachineInstr &
I)
const {
5865 MachineIRBuilder MIRBuilder(
I);
5868 ResType, MIRBuilder, SPIRV::StorageClass::Function);
5880 MachineBasicBlock &EntryBB =
I.getMF()->front();
5884 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
5887 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
5893 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5896 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
5899 .
add(
I.getOperand(
I.getNumExplicitDefs()))
5903 Register IntegralPartReg =
I.getOperand(1).getReg();
5904 if (IntegralPartReg.
isValid()) {
5906 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5917 assert(
false &&
"GLSL::Modf is deprecated.");
5928bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
5929 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5930 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5931 MachineIRBuilder MIRBuilder(
I);
5932 const SPIRVTypeInst Vec3Ty =
5935 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
5947 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5951 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
5957 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5964 assert(
I.getOperand(2).isReg());
5965 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
5969 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5980bool SPIRVInstructionSelector::loadBuiltinInputID(
5981 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5982 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5983 MachineIRBuilder MIRBuilder(
I);
5985 ResType, MIRBuilder, SPIRV::StorageClass::Input);
6000 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6004 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6013SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
6014 MachineInstr &
I)
const {
6015 MachineIRBuilder MIRBuilder(
I);
6016 if (
Type->getOpcode() != SPIRV::OpTypeVector)
6020 if (VectorSize == 4)
6028bool SPIRVInstructionSelector::loadHandleBeforePosition(
6029 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
6030 MachineInstr &Pos)
const {
6033 Intrinsic::spv_resource_handlefrombinding);
6041 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
6042 MachineIRBuilder MIRBuilder(HandleDef);
6043 SPIRVTypeInst VarType = ResType;
6044 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
6046 if (IsStructuredBuffer) {
6051 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
6053 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
6056 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
6057 ArraySize, IndexReg, Name, MIRBuilder);
6061 uint32_t LoadOpcode =
6062 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
6072void SPIRVInstructionSelector::errorIfInstrOutsideShader(
6073 MachineInstr &
I)
const {
6075 std::string DiagMsg;
6076 raw_string_ostream OS(DiagMsg);
6077 I.print(OS,
true,
false,
false,
false);
6078 DiagMsg +=
" is only supported in shaders.\n";
6084InstructionSelector *
6088 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 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.
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.
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)
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)
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)
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)
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...