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_named_boolean_spec_constant: {
4129 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4130 : SPIRV::OpSpecConstantFalse;
4132 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4133 .
addDef(
I.getOperand(0).getReg())
4136 unsigned SpecId =
I.getOperand(2).getImm();
4138 SPIRV::Decoration::SpecId, {SpecId});
4142 case Intrinsic::spv_const_composite: {
4144 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4150 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4152 MachineIRBuilder MIR(
I);
4154 MIR, SPIRV::OpConstantComposite, 3,
4155 SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,
4157 for (
auto *Instr : Instructions) {
4158 Instr->setDebugLoc(
I.getDebugLoc());
4163 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4170 case Intrinsic::spv_assign_name: {
4171 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4172 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4173 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4174 i <
I.getNumExplicitOperands(); ++i) {
4175 MIB.
addImm(
I.getOperand(i).getImm());
4180 case Intrinsic::spv_switch: {
4181 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4182 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4183 if (
I.getOperand(i).isReg())
4184 MIB.
addReg(
I.getOperand(i).getReg());
4185 else if (
I.getOperand(i).isCImm())
4186 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4187 else if (
I.getOperand(i).isMBB())
4188 MIB.
addMBB(
I.getOperand(i).getMBB());
4195 case Intrinsic::spv_loop_merge: {
4196 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4197 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4198 if (
I.getOperand(i).isMBB())
4199 MIB.
addMBB(
I.getOperand(i).getMBB());
4206 case Intrinsic::spv_loop_control_intel: {
4208 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4209 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4214 case Intrinsic::spv_selection_merge: {
4216 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4217 assert(
I.getOperand(1).isMBB() &&
4218 "operand 1 to spv_selection_merge must be a basic block");
4219 MIB.
addMBB(
I.getOperand(1).getMBB());
4220 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4224 case Intrinsic::spv_cmpxchg:
4225 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4226 case Intrinsic::spv_unreachable:
4227 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4230 case Intrinsic::spv_alloca:
4231 return selectFrameIndex(ResVReg, ResType,
I);
4232 case Intrinsic::spv_alloca_array:
4233 return selectAllocaArray(ResVReg, ResType,
I);
4234 case Intrinsic::spv_assume:
4236 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4237 .
addUse(
I.getOperand(1).getReg())
4242 case Intrinsic::spv_expect:
4244 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4247 .
addUse(
I.getOperand(2).getReg())
4248 .
addUse(
I.getOperand(3).getReg())
4253 case Intrinsic::arithmetic_fence:
4254 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4255 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4258 .
addUse(
I.getOperand(2).getReg())
4262 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4264 case Intrinsic::spv_thread_id:
4270 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4272 case Intrinsic::spv_thread_id_in_group:
4278 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4280 case Intrinsic::spv_group_id:
4286 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4288 case Intrinsic::spv_flattened_thread_id_in_group:
4295 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4297 case Intrinsic::spv_workgroup_size:
4298 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4300 case Intrinsic::spv_global_size:
4301 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4303 case Intrinsic::spv_global_offset:
4304 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4306 case Intrinsic::spv_num_workgroups:
4307 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4309 case Intrinsic::spv_subgroup_size:
4310 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4312 case Intrinsic::spv_num_subgroups:
4313 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4315 case Intrinsic::spv_subgroup_id:
4316 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4317 case Intrinsic::spv_subgroup_local_invocation_id:
4318 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4319 ResVReg, ResType,
I);
4320 case Intrinsic::spv_subgroup_max_size:
4321 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4323 case Intrinsic::spv_fdot:
4324 return selectFloatDot(ResVReg, ResType,
I);
4325 case Intrinsic::spv_udot:
4326 case Intrinsic::spv_sdot:
4327 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4329 return selectIntegerDot(ResVReg, ResType,
I,
4330 IID == Intrinsic::spv_sdot);
4331 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4332 case Intrinsic::spv_dot4add_i8packed:
4333 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4335 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4336 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4337 case Intrinsic::spv_dot4add_u8packed:
4338 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4340 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4341 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4342 case Intrinsic::spv_all:
4343 return selectAll(ResVReg, ResType,
I);
4344 case Intrinsic::spv_any:
4345 return selectAny(ResVReg, ResType,
I);
4346 case Intrinsic::spv_cross:
4347 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4348 case Intrinsic::spv_distance:
4349 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4350 case Intrinsic::spv_lerp:
4351 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4352 case Intrinsic::spv_length:
4353 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4354 case Intrinsic::spv_degrees:
4355 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4356 case Intrinsic::spv_faceforward:
4357 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4358 case Intrinsic::spv_frac:
4359 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4360 case Intrinsic::spv_isinf:
4361 return selectOpIsInf(ResVReg, ResType,
I);
4362 case Intrinsic::spv_isnan:
4363 return selectOpIsNan(ResVReg, ResType,
I);
4364 case Intrinsic::spv_normalize:
4365 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4366 case Intrinsic::spv_refract:
4367 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4368 case Intrinsic::spv_reflect:
4369 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4370 case Intrinsic::spv_rsqrt:
4371 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4372 case Intrinsic::spv_sign:
4373 return selectSign(ResVReg, ResType,
I);
4374 case Intrinsic::spv_smoothstep:
4375 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4376 case Intrinsic::spv_firstbituhigh:
4377 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4378 case Intrinsic::spv_firstbitshigh:
4379 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4380 case Intrinsic::spv_firstbitlow:
4381 return selectFirstBitLow(ResVReg, ResType,
I);
4382 case Intrinsic::spv_group_memory_barrier_with_group_sync: {
4384 buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent,
I);
4385 Register ScopeReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
4387 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpControlBarrier))
4394 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4395 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4396 SPIRV::StorageClass::StorageClass ResSC =
4400 "Generic storage class");
4401 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4409 case Intrinsic::spv_lifetime_start:
4410 case Intrinsic::spv_lifetime_end: {
4411 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
4412 : SPIRV::OpLifetimeStop;
4413 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
4414 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
4423 case Intrinsic::spv_saturate:
4424 return selectSaturate(ResVReg, ResType,
I);
4425 case Intrinsic::spv_nclamp:
4426 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
4427 case Intrinsic::spv_uclamp:
4428 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
4429 case Intrinsic::spv_sclamp:
4430 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
4431 case Intrinsic::spv_subgroup_prefix_bit_count:
4432 return selectWavePrefixBitCount(ResVReg, ResType,
I);
4433 case Intrinsic::spv_wave_active_countbits:
4434 return selectWaveActiveCountBits(ResVReg, ResType,
I);
4435 case Intrinsic::spv_wave_all_equal:
4436 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
4437 case Intrinsic::spv_wave_all:
4438 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
4439 case Intrinsic::spv_wave_any:
4440 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
4441 case Intrinsic::spv_subgroup_ballot:
4442 return selectWaveOpInst(ResVReg, ResType,
I,
4443 SPIRV::OpGroupNonUniformBallot);
4444 case Intrinsic::spv_wave_is_first_lane:
4445 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
4446 case Intrinsic::spv_wave_reduce_or:
4447 return selectWaveReduceOp(ResVReg, ResType,
I,
4448 SPIRV::OpGroupNonUniformBitwiseOr);
4449 case Intrinsic::spv_wave_reduce_xor:
4450 return selectWaveReduceOp(ResVReg, ResType,
I,
4451 SPIRV::OpGroupNonUniformBitwiseXor);
4452 case Intrinsic::spv_wave_reduce_umax:
4453 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
4454 case Intrinsic::spv_wave_reduce_max:
4455 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
4456 case Intrinsic::spv_wave_reduce_umin:
4457 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
4458 case Intrinsic::spv_wave_reduce_min:
4459 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
4460 case Intrinsic::spv_wave_reduce_sum:
4461 return selectWaveReduceSum(ResVReg, ResType,
I);
4462 case Intrinsic::spv_wave_product:
4463 return selectWaveReduceProduct(ResVReg, ResType,
I);
4464 case Intrinsic::spv_wave_readlane:
4465 return selectWaveOpInst(ResVReg, ResType,
I,
4466 SPIRV::OpGroupNonUniformShuffle);
4467 case Intrinsic::spv_wave_prefix_sum:
4468 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
4469 case Intrinsic::spv_wave_prefix_product:
4470 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
4471 case Intrinsic::spv_quad_read_across_x: {
4472 return selectQuadSwap(ResVReg, ResType,
I, 0);
4474 case Intrinsic::spv_step:
4475 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
4476 case Intrinsic::spv_radians:
4477 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
4481 case Intrinsic::instrprof_increment:
4482 case Intrinsic::instrprof_increment_step:
4483 case Intrinsic::instrprof_value_profile:
4486 case Intrinsic::spv_value_md:
4488 case Intrinsic::spv_resource_handlefrombinding: {
4489 return selectHandleFromBinding(ResVReg, ResType,
I);
4491 case Intrinsic::spv_resource_counterhandlefrombinding:
4492 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
4493 case Intrinsic::spv_resource_updatecounter:
4494 return selectUpdateCounter(ResVReg, ResType,
I);
4495 case Intrinsic::spv_resource_store_typedbuffer: {
4496 return selectImageWriteIntrinsic(
I);
4498 case Intrinsic::spv_resource_load_typedbuffer: {
4499 return selectReadImageIntrinsic(ResVReg, ResType,
I);
4501 case Intrinsic::spv_resource_load_level: {
4502 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
4504 case Intrinsic::spv_resource_sample:
4505 case Intrinsic::spv_resource_sample_clamp:
4506 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
4507 case Intrinsic::spv_resource_samplebias:
4508 case Intrinsic::spv_resource_samplebias_clamp:
4509 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
4510 case Intrinsic::spv_resource_samplegrad:
4511 case Intrinsic::spv_resource_samplegrad_clamp:
4512 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
4513 case Intrinsic::spv_resource_samplelevel:
4514 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
4515 case Intrinsic::spv_resource_samplecmp:
4516 case Intrinsic::spv_resource_samplecmp_clamp:
4517 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
4518 case Intrinsic::spv_resource_samplecmplevelzero:
4519 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
4520 case Intrinsic::spv_resource_gather:
4521 case Intrinsic::spv_resource_gather_cmp:
4522 return selectGatherIntrinsic(ResVReg, ResType,
I);
4523 case Intrinsic::spv_resource_getpointer: {
4524 return selectResourceGetPointer(ResVReg, ResType,
I);
4526 case Intrinsic::spv_pushconstant_getpointer: {
4527 return selectPushConstantGetPointer(ResVReg, ResType,
I);
4529 case Intrinsic::spv_discard: {
4530 return selectDiscard(ResVReg, ResType,
I);
4532 case Intrinsic::spv_resource_nonuniformindex: {
4533 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
4535 case Intrinsic::spv_unpackhalf2x16: {
4536 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
4538 case Intrinsic::spv_packhalf2x16: {
4539 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
4541 case Intrinsic::spv_ddx:
4542 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
4543 case Intrinsic::spv_ddy:
4544 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
4545 case Intrinsic::spv_ddx_coarse:
4546 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
4547 case Intrinsic::spv_ddy_coarse:
4548 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
4549 case Intrinsic::spv_ddx_fine:
4550 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
4551 case Intrinsic::spv_ddy_fine:
4552 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
4553 case Intrinsic::spv_fwidth:
4554 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
4555 case Intrinsic::spv_masked_gather:
4556 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4557 return selectMaskedGather(ResVReg, ResType,
I);
4558 return diagnoseUnsupported(
4559 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
4560 case Intrinsic::spv_masked_scatter:
4561 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4562 return selectMaskedScatter(
I);
4563 return diagnoseUnsupported(
4564 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
4566 std::string DiagMsg;
4567 raw_string_ostream OS(DiagMsg);
4569 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
4576bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
4577 SPIRVTypeInst ResType,
4578 MachineInstr &
I)
const {
4581 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
4588bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
4589 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4591 assert(Intr.getIntrinsicID() ==
4592 Intrinsic::spv_resource_counterhandlefrombinding);
4595 Register MainHandleReg = Intr.getOperand(2).getReg();
4597 assert(MainHandleDef->getIntrinsicID() ==
4598 Intrinsic::spv_resource_handlefrombinding);
4602 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
4603 Register IndexReg = MainHandleDef->getOperand(5).getReg();
4604 std::string CounterName =
4609 MachineIRBuilder MIRBuilder(
I);
4611 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
4613 ArraySize, IndexReg, CounterName, MIRBuilder);
4615 return BuildCOPY(ResVReg, CounterVarReg,
I);
4618bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
4619 SPIRVTypeInst ResType,
4620 MachineInstr &
I)
const {
4622 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
4624 Register CounterHandleReg = Intr.getOperand(2).getReg();
4625 Register IncrReg = Intr.getOperand(3).getReg();
4632 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
4633 assert(CounterVarPointeeType &&
4634 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
4635 "Counter variable must be a struct");
4637 SPIRV::StorageClass::StorageBuffer &&
4638 "Counter variable must be in the storage buffer storage class");
4640 "Counter variable must have exactly 1 member in the struct");
4641 const SPIRVTypeInst MemberType =
4644 "Counter variable struct must have a single i32 member");
4648 MachineIRBuilder MIRBuilder(
I);
4650 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
4653 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
4659 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4662 .
addUse(CounterHandleReg)
4669 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
4672 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
4675 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
4684 return BuildCOPY(ResVReg, AtomicRes,
I);
4692 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
4700bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
4701 SPIRVTypeInst ResType,
4702 MachineInstr &
I)
const {
4710 Register ImageReg =
I.getOperand(2).getReg();
4718 Register IdxReg =
I.getOperand(3).getReg();
4720 MachineInstr &Pos =
I;
4722 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
4726bool SPIRVInstructionSelector::generateSampleImage(
4729 DebugLoc Loc, MachineInstr &Pos)
const {
4740 if (!loadHandleBeforePosition(NewSamplerReg,
4746 MachineIRBuilder MIRBuilder(Pos);
4759 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
4760 ImOps.Lod.has_value();
4761 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
4762 : SPIRV::OpImageSampleImplicitLod;
4764 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
4765 : SPIRV::OpImageSampleDrefImplicitLod;
4774 MIB.
addUse(*ImOps.Compare);
4776 uint32_t ImageOperands = 0;
4778 ImageOperands |= SPIRV::ImageOperand::Bias;
4780 ImageOperands |= SPIRV::ImageOperand::Lod;
4781 if (ImOps.GradX && ImOps.GradY)
4782 ImageOperands |= SPIRV::ImageOperand::Grad;
4783 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
4785 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4788 "Non-constant offsets are not supported in sample instructions.");
4792 ImageOperands |= SPIRV::ImageOperand::MinLod;
4794 if (ImageOperands != 0) {
4795 MIB.
addImm(ImageOperands);
4796 if (ImageOperands & SPIRV::ImageOperand::Bias)
4798 if (ImageOperands & SPIRV::ImageOperand::Lod)
4800 if (ImageOperands & SPIRV::ImageOperand::Grad) {
4801 MIB.
addUse(*ImOps.GradX);
4802 MIB.
addUse(*ImOps.GradY);
4805 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4806 MIB.
addUse(*ImOps.Offset);
4807 if (ImageOperands & SPIRV::ImageOperand::MinLod)
4808 MIB.
addUse(*ImOps.MinLod);
4815bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
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 if (
I.getNumOperands() > 5)
4822 ImOps.Offset =
I.getOperand(5).getReg();
4823 if (
I.getNumOperands() > 6)
4824 ImOps.MinLod =
I.getOperand(6).getReg();
4825 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4826 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4829bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
4830 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4831 Register ImageReg =
I.getOperand(2).getReg();
4832 Register SamplerReg =
I.getOperand(3).getReg();
4833 Register CoordinateReg =
I.getOperand(4).getReg();
4834 ImageOperands ImOps;
4835 ImOps.Bias =
I.getOperand(5).getReg();
4836 if (
I.getNumOperands() > 6)
4837 ImOps.Offset =
I.getOperand(6).getReg();
4838 if (
I.getNumOperands() > 7)
4839 ImOps.MinLod =
I.getOperand(7).getReg();
4840 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4841 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4844bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
4845 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4846 Register ImageReg =
I.getOperand(2).getReg();
4847 Register SamplerReg =
I.getOperand(3).getReg();
4848 Register CoordinateReg =
I.getOperand(4).getReg();
4849 ImageOperands ImOps;
4850 ImOps.GradX =
I.getOperand(5).getReg();
4851 ImOps.GradY =
I.getOperand(6).getReg();
4852 if (
I.getNumOperands() > 7)
4853 ImOps.Offset =
I.getOperand(7).getReg();
4854 if (
I.getNumOperands() > 8)
4855 ImOps.MinLod =
I.getOperand(8).getReg();
4856 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4857 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4860bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
4861 Register &ResVReg, SPIRVTypeInst ResType, 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.Lod =
I.getOperand(5).getReg();
4867 if (
I.getNumOperands() > 6)
4868 ImOps.Offset =
I.getOperand(6).getReg();
4869 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4870 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4873bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
4874 SPIRVTypeInst ResType,
4875 MachineInstr &
I)
const {
4876 Register ImageReg =
I.getOperand(2).getReg();
4877 Register SamplerReg =
I.getOperand(3).getReg();
4878 Register CoordinateReg =
I.getOperand(4).getReg();
4879 ImageOperands ImOps;
4880 ImOps.Compare =
I.getOperand(5).getReg();
4881 if (
I.getNumOperands() > 6)
4882 ImOps.Offset =
I.getOperand(6).getReg();
4883 if (
I.getNumOperands() > 7)
4884 ImOps.MinLod =
I.getOperand(7).getReg();
4885 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4886 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4889bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
4890 SPIRVTypeInst ResType,
4891 MachineInstr &
I)
const {
4892 Register ImageReg =
I.getOperand(2).getReg();
4893 Register CoordinateReg =
I.getOperand(3).getReg();
4894 Register LodReg =
I.getOperand(4).getReg();
4896 ImageOperands ImOps;
4898 if (
I.getNumOperands() > 5)
4899 ImOps.Offset =
I.getOperand(5).getReg();
4911 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
4912 I.getDebugLoc(),
I, &ImOps);
4915bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
4916 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4917 Register ImageReg =
I.getOperand(2).getReg();
4918 Register SamplerReg =
I.getOperand(3).getReg();
4919 Register CoordinateReg =
I.getOperand(4).getReg();
4920 ImageOperands ImOps;
4921 ImOps.Compare =
I.getOperand(5).getReg();
4922 if (
I.getNumOperands() > 6)
4923 ImOps.Offset =
I.getOperand(6).getReg();
4926 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4927 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4930bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
4931 SPIRVTypeInst ResType,
4932 MachineInstr &
I)
const {
4933 Register ImageReg =
I.getOperand(2).getReg();
4934 Register SamplerReg =
I.getOperand(3).getReg();
4935 Register CoordinateReg =
I.getOperand(4).getReg();
4938 "ImageReg is not an image type.");
4943 ComponentOrCompareReg =
I.getOperand(5).getReg();
4944 OffsetReg =
I.getOperand(6).getReg();
4947 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
4951 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
4952 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
4953 Dim != SPIRV::Dim::DIM_Rect) {
4955 "Gather operations are only supported for 2D, Cube, and Rect images.");
4962 if (!loadHandleBeforePosition(
4967 MachineIRBuilder MIRBuilder(
I);
4968 SPIRVTypeInst SampledImageType =
4973 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
4981 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
4983 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
4985 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4990 .
addUse(ComponentOrCompareReg);
4992 uint32_t ImageOperands = 0;
4993 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
4994 if (Dim == SPIRV::Dim::DIM_Cube) {
4996 "Gather operations with offset are not supported for Cube images.");
5000 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5002 ImageOperands |= SPIRV::ImageOperand::Offset;
5006 if (ImageOperands != 0) {
5007 MIB.
addImm(ImageOperands);
5009 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5017bool SPIRVInstructionSelector::generateImageReadOrFetch(
5020 const ImageOperands *ImOps)
const {
5023 "ImageReg is not an image type.");
5025 bool IsSignedInteger =
5030 bool IsFetch = (SampledOp.getImm() == 1);
5032 auto AddOperands = [&](MachineInstrBuilder &MIB) {
5033 uint32_t ImageOperandsMask = 0;
5034 if (IsSignedInteger)
5035 ImageOperandsMask |= 0x1000;
5037 if (IsFetch && ImOps) {
5039 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
5040 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
5042 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
5044 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5048 if (ImageOperandsMask != 0) {
5049 MIB.
addImm(ImageOperandsMask);
5050 if (IsFetch && ImOps) {
5053 if (ImOps->Offset &&
5054 (ImageOperandsMask &
5055 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5056 MIB.
addUse(*ImOps->Offset);
5062 if (ResultSize == 4) {
5065 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5072 BMI.constrainAllUses(
TII,
TRI, RBI);
5076 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5080 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5086 BMI.constrainAllUses(
TII,
TRI, RBI);
5088 if (ResultSize == 1) {
5097 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
5100bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
5101 SPIRVTypeInst ResType,
5102 MachineInstr &
I)
const {
5103 Register ResourcePtr =
I.getOperand(2).getReg();
5105 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
5114 MachineIRBuilder MIRBuilder(
I);
5116 Register IndexReg =
I.getOperand(3).getReg();
5119 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5129bool SPIRVInstructionSelector::selectPushConstantGetPointer(
5130 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5135bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
5136 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5137 Register ObjReg =
I.getOperand(2).getReg();
5138 if (!BuildCOPY(ResVReg, ObjReg,
I))
5148 decorateUsesAsNonUniform(ResVReg);
5152void SPIRVInstructionSelector::decorateUsesAsNonUniform(
5155 while (WorkList.
size() > 0) {
5159 bool IsDecorated =
false;
5161 if (
Use.getOpcode() == SPIRV::OpDecorate &&
5162 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
5168 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
5170 if (ResultReg == CurrentReg)
5178 SPIRV::Decoration::NonUniformEXT, {});
5183bool SPIRVInstructionSelector::extractSubvector(
5185 MachineInstr &InsertionPoint)
const {
5187 [[maybe_unused]] uint64_t InputSize =
5190 assert(InputSize > 1 &&
"The input must be a vector.");
5191 assert(ResultSize > 1 &&
"The result must be a vector.");
5192 assert(ResultSize < InputSize &&
5193 "Cannot extract more element than there are in the input.");
5196 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
5197 for (uint64_t
I = 0;
I < ResultSize;
I++) {
5200 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5209 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
5211 TII.get(SPIRV::OpCompositeConstruct))
5215 for (
Register ComponentReg : ComponentRegisters)
5216 MIB.
addUse(ComponentReg);
5221bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
5222 MachineInstr &
I)
const {
5229 Register ImageReg =
I.getOperand(1).getReg();
5237 Register CoordinateReg =
I.getOperand(2).getReg();
5238 Register DataReg =
I.getOperand(3).getReg();
5241 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
5249Register SPIRVInstructionSelector::buildPointerToResource(
5250 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
5251 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
5252 StringRef Name, MachineIRBuilder MIRBuilder)
const {
5254 if (ArraySize == 1) {
5255 SPIRVTypeInst PtrType =
5258 "SpirvResType did not have an explicit layout.");
5263 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
5264 SPIRVTypeInst VarPointerType =
5267 VarPointerType, Set,
Binding, Name, MIRBuilder);
5269 SPIRVTypeInst ResPointerType =
5282bool SPIRVInstructionSelector::selectFirstBitSet16(
5283 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5284 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
5286 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
5290 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
5293bool SPIRVInstructionSelector::selectFirstBitSet32(
5295 unsigned BitSetOpcode)
const {
5296 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5299 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5306bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
5308 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5315 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
5317 MachineIRBuilder MIRBuilder(
I);
5320 SPIRVTypeInst I64x2Type =
5322 SPIRVTypeInst Vec2ResType =
5325 std::vector<Register> PartialRegs;
5328 unsigned CurrentComponent = 0;
5329 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
5335 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5336 TII.get(SPIRV::OpVectorShuffle))
5341 .
addImm(CurrentComponent)
5342 .
addImm(CurrentComponent + 1);
5349 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType,
I, BitSetResult,
5350 BitSetOpcode, SwapPrimarySide))
5353 PartialRegs.push_back(SubVecBitSetReg);
5357 if (CurrentComponent != ComponentCount) {
5363 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
5364 SPIRV::OpVectorExtractDynamic))
5370 if (!selectFirstBitSet64(FinalElemBitSetReg,
BaseType,
I, FinalElemReg,
5371 BitSetOpcode, SwapPrimarySide))
5374 PartialRegs.push_back(FinalElemBitSetReg);
5379 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
5380 SPIRV::OpCompositeConstruct);
5383bool SPIRVInstructionSelector::selectFirstBitSet64(
5385 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5398 if (ComponentCount > 2) {
5399 return selectFirstBitSet64Overflow(ResVReg, ResType,
I, SrcReg,
5400 BitSetOpcode, SwapPrimarySide);
5404 MachineIRBuilder MIRBuilder(
I);
5406 BaseType, 2 * ComponentCount, MIRBuilder,
false);
5410 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
5416 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
5423 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
5426 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntZero},
5427 SPIRV::OpVectorExtractDynamic))
5429 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntOne},
5430 SPIRV::OpVectorExtractDynamic))
5434 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5435 TII.get(SPIRV::OpVectorShuffle))
5443 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
5449 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5450 TII.get(SPIRV::OpVectorShuffle))
5458 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
5478 SelectOp = SPIRV::OpSelectSISCond;
5479 AddOp = SPIRV::OpIAddS;
5487 SelectOp = SPIRV::OpSelectVIVCond;
5488 AddOp = SPIRV::OpIAddV;
5498 if (SwapPrimarySide) {
5499 PrimaryReg = LowReg;
5500 SecondaryReg = HighReg;
5501 PrimaryShiftReg = Reg0;
5502 SecondaryShiftReg = Reg32;
5507 if (!selectOpWithSrcs(BReg, BoolType,
I, {PrimaryReg, NegOneReg},
5513 if (!selectOpWithSrcs(TmpReg, ResType,
I, {
BReg, SecondaryReg, PrimaryReg},
5519 if (!selectOpWithSrcs(ValReg, ResType,
I,
5520 {
BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))
5523 return selectOpWithSrcs(ResVReg, ResType,
I, {ValReg, TmpReg}, AddOp);
5526bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
5527 SPIRVTypeInst ResType,
5529 bool IsSigned)
const {
5531 Register OpReg =
I.getOperand(2).getReg();
5534 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5535 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
5539 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5541 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5543 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5547 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
5551bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
5552 SPIRVTypeInst ResType,
5553 MachineInstr &
I)
const {
5555 Register OpReg =
I.getOperand(2).getReg();
5560 unsigned ExtendOpcode = SPIRV::OpUConvert;
5561 unsigned BitSetOpcode = GL::FindILsb;
5565 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5567 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5569 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5576bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
5577 SPIRVTypeInst ResType,
5578 MachineInstr &
I)
const {
5582 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
5585 .
addUse(
I.getOperand(2).getReg())
5588 unsigned Alignment =
I.getOperand(3).getImm();
5594bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
5595 SPIRVTypeInst ResType,
5596 MachineInstr &
I)
const {
5600 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
5603 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
5606 unsigned Alignment =
I.getOperand(2).getImm();
5613bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
5618 const MachineInstr *PrevI =
I.getPrevNode();
5620 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
5624 .
addMBB(
I.getOperand(0).getMBB())
5629 .
addMBB(
I.getOperand(0).getMBB())
5634bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
5645 const MachineInstr *NextI =
I.getNextNode();
5647 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
5653 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
5655 .
addUse(
I.getOperand(0).getReg())
5656 .
addMBB(
I.getOperand(1).getMBB())
5662bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
5663 MachineInstr &
I)
const {
5665 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
5667 const unsigned NumOps =
I.getNumOperands();
5668 for (
unsigned i = 1; i <
NumOps; i += 2) {
5669 MIB.
addUse(
I.getOperand(i + 0).getReg());
5670 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
5676bool SPIRVInstructionSelector::selectGlobalValue(
5677 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
5679 MachineIRBuilder MIRBuilder(
I);
5680 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
5683 std::string GlobalIdent;
5685 unsigned &
ID = UnnamedGlobalIDs[GV];
5687 ID = UnnamedGlobalIDs.
size();
5688 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
5714 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
5721 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
5726 MachineInstrBuilder MIB1 =
5727 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5730 MachineInstrBuilder MIB2 =
5732 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
5736 GR.
add(ConstVal, MIB2);
5744 MachineInstrBuilder MIB3 =
5745 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5748 GR.
add(ConstVal, MIB3);
5752 assert(NewReg != ResVReg);
5753 return BuildCOPY(ResVReg, NewReg,
I);
5763 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
5769 SPIRVTypeInst ResType =
5773 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
5778 if (
GlobalVar->isExternallyInitialized() &&
5779 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
5780 constexpr unsigned ReadWriteINTEL = 3u;
5783 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
5789bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
5790 SPIRVTypeInst ResType,
5791 MachineInstr &
I)
const {
5793 return selectExtInst(ResVReg, ResType,
I, CL::log10);
5801 MachineIRBuilder MIRBuilder(
I);
5806 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5809 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5811 .
add(
I.getOperand(1))
5816 ResType->
getOpcode() == SPIRV::OpTypeFloat);
5818 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
5826 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
5827 ? SPIRV::OpVectorTimesScalar
5838bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
5839 SPIRVTypeInst ResType,
5840 MachineInstr &
I)
const {
5843 return selectExtInst(ResVReg, ResType,
I, CL::pown);
5849 Register ExpReg =
I.getOperand(2).getReg();
5851 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
5852 SPIRV::OpConvertSToF))
5854 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
5861bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
5862 SPIRVTypeInst ResType,
5863 MachineInstr &
I)
const {
5879 MachineIRBuilder MIRBuilder(
I);
5882 ResType, MIRBuilder, SPIRV::StorageClass::Function);
5894 MachineBasicBlock &EntryBB =
I.getMF()->front();
5898 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
5901 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
5907 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5910 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
5913 .
add(
I.getOperand(
I.getNumExplicitDefs()))
5917 Register IntegralPartReg =
I.getOperand(1).getReg();
5918 if (IntegralPartReg.
isValid()) {
5920 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5931 assert(
false &&
"GLSL::Modf is deprecated.");
5942bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
5943 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5944 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5945 MachineIRBuilder MIRBuilder(
I);
5946 const SPIRVTypeInst Vec3Ty =
5949 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
5961 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5965 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
5971 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5978 assert(
I.getOperand(2).isReg());
5979 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
5983 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5994bool SPIRVInstructionSelector::loadBuiltinInputID(
5995 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5996 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5997 MachineIRBuilder MIRBuilder(
I);
5999 ResType, MIRBuilder, SPIRV::StorageClass::Input);
6014 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6018 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6027SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
6028 MachineInstr &
I)
const {
6029 MachineIRBuilder MIRBuilder(
I);
6030 if (
Type->getOpcode() != SPIRV::OpTypeVector)
6034 if (VectorSize == 4)
6042bool SPIRVInstructionSelector::loadHandleBeforePosition(
6043 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
6044 MachineInstr &Pos)
const {
6047 Intrinsic::spv_resource_handlefrombinding);
6055 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
6056 MachineIRBuilder MIRBuilder(HandleDef);
6057 SPIRVTypeInst VarType = ResType;
6058 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
6060 if (IsStructuredBuffer) {
6065 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
6067 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
6070 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
6071 ArraySize, IndexReg, Name, MIRBuilder);
6075 uint32_t LoadOpcode =
6076 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
6086void SPIRVInstructionSelector::errorIfInstrOutsideShader(
6087 MachineInstr &
I)
const {
6089 std::string DiagMsg;
6090 raw_string_ostream OS(DiagMsg);
6091 I.print(OS,
true,
false,
false,
false);
6092 DiagMsg +=
" is only supported in shaders.\n";
6098InstructionSelector *
6102 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...