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;
230 template <
bool Signed>
233 template <
bool Signed>
240 template <
typename PickOpcodeFn>
243 PickOpcodeFn &&PickOpcode)
const;
260 template <
typename PickOpcodeFn>
263 PickOpcodeFn &&PickOpcode)
const;
278 bool IsSigned)
const;
280 bool IsSigned,
unsigned Opcode)
const;
282 bool IsSigned)
const;
288 bool IsSigned)
const;
327 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
328 bool useMISrc =
true,
330 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
331 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
332 bool useMISrc =
true,
334 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
335 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
336 bool setMIFlags =
true,
bool useMISrc =
true,
338 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
339 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
340 bool useMISrc =
true,
343 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
344 MachineInstr &
I)
const;
346 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
347 MachineInstr &
I)
const;
349 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
350 MachineInstr &
I)
const;
352 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
353 MachineInstr &
I,
unsigned Opcode)
const;
355 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
356 MachineInstr &
I)
const;
358 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
359 MachineInstr &
I)
const;
363 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
364 MachineInstr &
I)
const;
366 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
367 MachineInstr &
I)
const;
369 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
370 MachineInstr &
I)
const;
371 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
372 MachineInstr &
I)
const;
373 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
374 MachineInstr &
I)
const;
375 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
376 MachineInstr &
I)
const;
377 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
378 MachineInstr &
I)
const;
379 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
380 MachineInstr &
I)
const;
381 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
382 SPIRVTypeInst ResType,
383 MachineInstr &
I)
const;
384 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
385 MachineInstr &
I)
const;
386 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
387 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
388 MachineInstr &
I)
const;
389 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
390 MachineInstr &
I)
const;
391 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
392 MachineInstr &
I)
const;
393 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
394 MachineInstr &
I)
const;
395 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
396 MachineInstr &
I)
const;
397 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
398 MachineInstr &
I)
const;
399 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
400 MachineInstr &
I)
const;
401 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
402 MachineInstr &
I)
const;
403 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
404 MachineInstr &
I,
const unsigned DPdOpCode)
const;
406 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
407 SPIRVTypeInst ResType =
nullptr)
const;
409 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
410 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
411 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
413 MachineInstr &
I)
const;
414 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
416 bool wrapIntoSpecConstantOp(MachineInstr &
I,
419 Register getUcharPtrTypeReg(MachineInstr &
I,
420 SPIRV::StorageClass::StorageClass SC)
const;
421 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
423 uint32_t Opcode)
const;
424 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
425 SPIRVTypeInst SrcPtrTy)
const;
426 Register buildPointerToResource(SPIRVTypeInst ResType,
427 SPIRV::StorageClass::StorageClass SC,
428 uint32_t Set, uint32_t
Binding,
429 uint32_t ArraySize,
Register IndexReg,
431 MachineIRBuilder MIRBuilder)
const;
432 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
433 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
434 Register &ReadReg, MachineInstr &InsertionPoint)
const;
435 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
437 DebugLoc Loc, MachineInstr &Pos)
const;
438 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
440 Register CoordinateReg,
const ImageOperands &ImOps,
443 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
444 Register ResVReg, SPIRVTypeInst ResType,
445 MachineInstr &
I)
const;
446 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
447 Register ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
450 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
451 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
452 void errorIfInstrOutsideShader(MachineInstr &
I)
const;
455bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
457 if (
TET->getTargetExtName() ==
"spirv.Image") {
460 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
461 return TET->getTypeParameter(0)->isIntegerTy();
465#define GET_GLOBALISEL_IMPL
466#include "SPIRVGenGlobalISel.inc"
467#undef GET_GLOBALISEL_IMPL
473 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
476#include
"SPIRVGenGlobalISel.inc"
479#include
"SPIRVGenGlobalISel.inc"
491 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
495void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
496 if (HasVRegsReset == &MF)
511 for (
const auto &
MBB : MF) {
512 for (
const auto &
MI :
MBB) {
515 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
519 LLT DstType = MRI.
getType(DstReg);
521 LLT SrcType = MRI.
getType(SrcReg);
522 if (DstType != SrcType)
527 if (DstRC != SrcRC && SrcRC)
539 while (!Stack.empty()) {
544 switch (
MI->getOpcode()) {
545 case TargetOpcode::G_INTRINSIC:
546 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
547 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
549 Intrinsic::spv_const_composite)
552 case TargetOpcode::G_BUILD_VECTOR:
553 case TargetOpcode::G_SPLAT_VECTOR:
555 i < OpDef->getNumOperands(); i++) {
560 Stack.push_back(OpNestedDef);
563 case TargetOpcode::G_CONSTANT:
564 case TargetOpcode::G_FCONSTANT:
565 case TargetOpcode::G_IMPLICIT_DEF:
566 case SPIRV::OpConstantTrue:
567 case SPIRV::OpConstantFalse:
568 case SPIRV::OpConstantI:
569 case SPIRV::OpConstantF:
570 case SPIRV::OpConstantComposite:
571 case SPIRV::OpConstantCompositeContinuedINTEL:
572 case SPIRV::OpConstantSampler:
573 case SPIRV::OpConstantNull:
575 case SPIRV::OpConstantFunctionPointerINTEL:
602 case Intrinsic::spv_all:
603 case Intrinsic::spv_alloca:
604 case Intrinsic::spv_any:
605 case Intrinsic::spv_bitcast:
606 case Intrinsic::spv_const_composite:
607 case Intrinsic::spv_cross:
608 case Intrinsic::spv_degrees:
609 case Intrinsic::spv_distance:
610 case Intrinsic::spv_extractelt:
611 case Intrinsic::spv_extractv:
612 case Intrinsic::spv_faceforward:
613 case Intrinsic::spv_fdot:
614 case Intrinsic::spv_firstbitlow:
615 case Intrinsic::spv_firstbitshigh:
616 case Intrinsic::spv_firstbituhigh:
617 case Intrinsic::spv_frac:
618 case Intrinsic::spv_gep:
619 case Intrinsic::spv_global_offset:
620 case Intrinsic::spv_global_size:
621 case Intrinsic::spv_group_id:
622 case Intrinsic::spv_insertelt:
623 case Intrinsic::spv_insertv:
624 case Intrinsic::spv_isinf:
625 case Intrinsic::spv_isnan:
626 case Intrinsic::spv_lerp:
627 case Intrinsic::spv_length:
628 case Intrinsic::spv_normalize:
629 case Intrinsic::spv_num_subgroups:
630 case Intrinsic::spv_num_workgroups:
631 case Intrinsic::spv_ptrcast:
632 case Intrinsic::spv_radians:
633 case Intrinsic::spv_reflect:
634 case Intrinsic::spv_refract:
635 case Intrinsic::spv_resource_getpointer:
636 case Intrinsic::spv_resource_handlefrombinding:
637 case Intrinsic::spv_resource_handlefromimplicitbinding:
638 case Intrinsic::spv_resource_nonuniformindex:
639 case Intrinsic::spv_resource_sample:
640 case Intrinsic::spv_rsqrt:
641 case Intrinsic::spv_saturate:
642 case Intrinsic::spv_sdot:
643 case Intrinsic::spv_sign:
644 case Intrinsic::spv_smoothstep:
645 case Intrinsic::spv_step:
646 case Intrinsic::spv_subgroup_id:
647 case Intrinsic::spv_subgroup_local_invocation_id:
648 case Intrinsic::spv_subgroup_max_size:
649 case Intrinsic::spv_subgroup_size:
650 case Intrinsic::spv_thread_id:
651 case Intrinsic::spv_thread_id_in_group:
652 case Intrinsic::spv_udot:
653 case Intrinsic::spv_undef:
654 case Intrinsic::spv_value_md:
655 case Intrinsic::spv_workgroup_size:
667 case SPIRV::OpTypeVoid:
668 case SPIRV::OpTypeBool:
669 case SPIRV::OpTypeInt:
670 case SPIRV::OpTypeFloat:
671 case SPIRV::OpTypeVector:
672 case SPIRV::OpTypeMatrix:
673 case SPIRV::OpTypeImage:
674 case SPIRV::OpTypeSampler:
675 case SPIRV::OpTypeSampledImage:
676 case SPIRV::OpTypeArray:
677 case SPIRV::OpTypeRuntimeArray:
678 case SPIRV::OpTypeStruct:
679 case SPIRV::OpTypeOpaque:
680 case SPIRV::OpTypePointer:
681 case SPIRV::OpTypeFunction:
682 case SPIRV::OpTypeEvent:
683 case SPIRV::OpTypeDeviceEvent:
684 case SPIRV::OpTypeReserveId:
685 case SPIRV::OpTypeQueue:
686 case SPIRV::OpTypePipe:
687 case SPIRV::OpTypeForwardPointer:
688 case SPIRV::OpTypePipeStorage:
689 case SPIRV::OpTypeNamedBarrier:
690 case SPIRV::OpTypeAccelerationStructureNV:
691 case SPIRV::OpTypeCooperativeMatrixNV:
692 case SPIRV::OpTypeCooperativeMatrixKHR:
702 if (
MI.getNumDefs() == 0)
705 for (
const auto &MO :
MI.all_defs()) {
707 if (
Reg.isPhysical()) {
712 if (
UseMI.getOpcode() != SPIRV::OpName) {
719 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
720 MI.isLifetimeMarker()) {
723 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
734 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
735 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
738 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
743 if (
MI.mayStore() ||
MI.isCall() ||
744 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
745 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
746 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
757 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
764void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
766 for (
const auto &MO :
MI.all_defs()) {
770 SmallVector<MachineInstr *, 4> UselessOpNames;
773 "There is still a use of the dead function.");
776 for (MachineInstr *OpNameMI : UselessOpNames) {
778 OpNameMI->eraseFromParent();
783void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
786 removeOpNamesForDeadMI(
MI);
787 MI.eraseFromParent();
790bool SPIRVInstructionSelector::select(MachineInstr &
I) {
791 resetVRegsType(*
I.getParent()->getParent());
793 assert(
I.getParent() &&
"Instruction should be in a basic block!");
794 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
799 removeDeadInstruction(
I);
806 if (Opcode == SPIRV::ASSIGN_TYPE) {
807 Register DstReg =
I.getOperand(0).getReg();
808 Register SrcReg =
I.getOperand(1).getReg();
811 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
812 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
813 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
814 Register SelectDstReg =
Def->getOperand(0).getReg();
815 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
817 assert(SuccessToSelectSelect);
819 Def->eraseFromParent();
826 bool Res = selectImpl(
I, *CoverageInfo);
828 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
829 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
833 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
845 }
else if (
I.getNumDefs() == 1) {
857 removeDeadInstruction(
I);
862 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
863 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
869 bool HasDefs =
I.getNumDefs() > 0;
872 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
873 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
874 if (spvSelect(ResVReg, ResType,
I)) {
876 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
887 case TargetOpcode::G_CONSTANT:
888 case TargetOpcode::G_FCONSTANT:
890 case TargetOpcode::G_SADDO:
891 case TargetOpcode::G_SSUBO:
898 MachineInstr &
I)
const {
901 if (DstRC != SrcRC && SrcRC)
903 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
910bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
911 SPIRVTypeInst ResType,
912 MachineInstr &
I)
const {
913 const unsigned Opcode =
I.getOpcode();
915 return selectImpl(
I, *CoverageInfo);
917 case TargetOpcode::G_CONSTANT:
918 case TargetOpcode::G_FCONSTANT:
919 return selectConst(ResVReg, ResType,
I);
920 case TargetOpcode::G_GLOBAL_VALUE:
921 return selectGlobalValue(ResVReg,
I);
922 case TargetOpcode::G_IMPLICIT_DEF:
923 return selectOpUndef(ResVReg, ResType,
I);
924 case TargetOpcode::G_FREEZE:
925 return selectFreeze(ResVReg, ResType,
I);
927 case TargetOpcode::G_INTRINSIC:
928 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
929 case TargetOpcode::G_INTRINSIC_CONVERGENT:
930 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
931 return selectIntrinsic(ResVReg, ResType,
I);
932 case TargetOpcode::G_BITREVERSE:
933 return selectBitreverse(ResVReg, ResType,
I);
935 case TargetOpcode::G_BUILD_VECTOR:
936 return selectBuildVector(ResVReg, ResType,
I);
937 case TargetOpcode::G_SPLAT_VECTOR:
938 return selectSplatVector(ResVReg, ResType,
I);
940 case TargetOpcode::G_SHUFFLE_VECTOR: {
941 MachineBasicBlock &BB = *
I.getParent();
942 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
945 .
addUse(
I.getOperand(1).getReg())
946 .
addUse(
I.getOperand(2).getReg());
947 for (
auto V :
I.getOperand(3).getShuffleMask())
952 case TargetOpcode::G_MEMMOVE:
953 case TargetOpcode::G_MEMCPY:
954 case TargetOpcode::G_MEMSET:
955 return selectMemOperation(ResVReg,
I);
957 case TargetOpcode::G_ICMP:
958 return selectICmp(ResVReg, ResType,
I);
959 case TargetOpcode::G_FCMP:
960 return selectFCmp(ResVReg, ResType,
I);
962 case TargetOpcode::G_FRAME_INDEX:
963 return selectFrameIndex(ResVReg, ResType,
I);
965 case TargetOpcode::G_LOAD:
966 return selectLoad(ResVReg, ResType,
I);
967 case TargetOpcode::G_STORE:
968 return selectStore(
I);
970 case TargetOpcode::G_BR:
971 return selectBranch(
I);
972 case TargetOpcode::G_BRCOND:
973 return selectBranchCond(
I);
975 case TargetOpcode::G_PHI:
976 return selectPhi(ResVReg,
I);
978 case TargetOpcode::G_FPTOSI:
979 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
980 case TargetOpcode::G_FPTOUI:
981 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
983 case TargetOpcode::G_FPTOSI_SAT:
984 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
985 case TargetOpcode::G_FPTOUI_SAT:
986 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
988 case TargetOpcode::G_SITOFP:
989 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
990 case TargetOpcode::G_UITOFP:
991 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
993 case TargetOpcode::G_CTPOP:
994 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitCount);
995 case TargetOpcode::G_SMIN:
996 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
997 case TargetOpcode::G_UMIN:
998 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1000 case TargetOpcode::G_SMAX:
1001 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1002 case TargetOpcode::G_UMAX:
1003 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1005 case TargetOpcode::G_SCMP:
1006 return selectSUCmp(ResVReg, ResType,
I,
true);
1007 case TargetOpcode::G_UCMP:
1008 return selectSUCmp(ResVReg, ResType,
I,
false);
1009 case TargetOpcode::G_LROUND:
1010 case TargetOpcode::G_LLROUND: {
1013 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1015 regForLround, *(
I.getParent()->getParent()));
1017 CL::round, GL::Round,
false);
1019 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1026 case TargetOpcode::G_STRICT_FMA:
1027 case TargetOpcode::G_FMA: {
1030 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1033 .
addUse(
I.getOperand(1).getReg())
1034 .
addUse(
I.getOperand(2).getReg())
1035 .
addUse(
I.getOperand(3).getReg())
1040 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1043 case TargetOpcode::G_STRICT_FLDEXP:
1044 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1046 case TargetOpcode::G_FPOW:
1047 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1048 case TargetOpcode::G_FPOWI:
1049 return selectFpowi(ResVReg, ResType,
I);
1051 case TargetOpcode::G_FEXP:
1052 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1053 case TargetOpcode::G_FEXP2:
1054 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1055 case TargetOpcode::G_FEXP10:
1056 return selectExp10(ResVReg, ResType,
I);
1058 case TargetOpcode::G_FMODF:
1059 return selectModf(ResVReg, ResType,
I);
1060 case TargetOpcode::G_FSINCOS:
1061 return selectSincos(ResVReg, ResType,
I);
1063 case TargetOpcode::G_FLOG:
1064 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1065 case TargetOpcode::G_FLOG2:
1066 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1067 case TargetOpcode::G_FLOG10:
1068 return selectLog10(ResVReg, ResType,
I);
1070 case TargetOpcode::G_FABS:
1071 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1072 case TargetOpcode::G_ABS:
1073 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1075 case TargetOpcode::G_FMINNUM:
1076 case TargetOpcode::G_FMINIMUM:
1077 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1078 case TargetOpcode::G_FMAXNUM:
1079 case TargetOpcode::G_FMAXIMUM:
1080 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1082 case TargetOpcode::G_FCOPYSIGN:
1083 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1085 case TargetOpcode::G_FCEIL:
1086 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1087 case TargetOpcode::G_FFLOOR:
1088 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1090 case TargetOpcode::G_FCOS:
1091 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1092 case TargetOpcode::G_FSIN:
1093 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1094 case TargetOpcode::G_FTAN:
1095 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1096 case TargetOpcode::G_FACOS:
1097 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1098 case TargetOpcode::G_FASIN:
1099 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1100 case TargetOpcode::G_FATAN:
1101 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1102 case TargetOpcode::G_FATAN2:
1103 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1104 case TargetOpcode::G_FCOSH:
1105 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1106 case TargetOpcode::G_FSINH:
1107 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1108 case TargetOpcode::G_FTANH:
1109 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1111 case TargetOpcode::G_STRICT_FSQRT:
1112 case TargetOpcode::G_FSQRT:
1113 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1115 case TargetOpcode::G_CTTZ:
1116 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1117 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1118 case TargetOpcode::G_CTLZ:
1119 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1120 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1122 case TargetOpcode::G_INTRINSIC_ROUND:
1123 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1124 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1125 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1126 case TargetOpcode::G_INTRINSIC_TRUNC:
1127 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1128 case TargetOpcode::G_FRINT:
1129 case TargetOpcode::G_FNEARBYINT:
1130 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1132 case TargetOpcode::G_SMULH:
1133 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1134 case TargetOpcode::G_UMULH:
1135 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1137 case TargetOpcode::G_SADDSAT:
1138 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1139 case TargetOpcode::G_UADDSAT:
1140 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1141 case TargetOpcode::G_SSUBSAT:
1142 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1143 case TargetOpcode::G_USUBSAT:
1144 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1146 case TargetOpcode::G_FFREXP:
1147 return selectFrexp(ResVReg, ResType,
I);
1149 case TargetOpcode::G_UADDO:
1150 return selectOverflowArith(ResVReg, ResType,
I,
1151 ResType->
getOpcode() == SPIRV::OpTypeVector
1152 ? SPIRV::OpIAddCarryV
1153 : SPIRV::OpIAddCarryS);
1154 case TargetOpcode::G_USUBO:
1155 return selectOverflowArith(ResVReg, ResType,
I,
1156 ResType->
getOpcode() == SPIRV::OpTypeVector
1157 ? SPIRV::OpISubBorrowV
1158 : SPIRV::OpISubBorrowS);
1159 case TargetOpcode::G_UMULO:
1160 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1161 case TargetOpcode::G_SMULO:
1162 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1164 case TargetOpcode::G_SEXT:
1165 return selectExt(ResVReg, ResType,
I,
true);
1166 case TargetOpcode::G_ANYEXT:
1167 case TargetOpcode::G_ZEXT:
1168 return selectExt(ResVReg, ResType,
I,
false);
1169 case TargetOpcode::G_TRUNC:
1170 return selectTrunc(ResVReg, ResType,
I);
1171 case TargetOpcode::G_FPTRUNC:
1172 case TargetOpcode::G_FPEXT:
1173 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1175 case TargetOpcode::G_PTRTOINT:
1176 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1177 case TargetOpcode::G_INTTOPTR:
1178 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1179 case TargetOpcode::G_BITCAST:
1180 return selectBitcast(ResVReg, ResType,
I);
1181 case TargetOpcode::G_ADDRSPACE_CAST:
1182 return selectAddrSpaceCast(ResVReg, ResType,
I);
1183 case TargetOpcode::G_PTR_ADD: {
1185 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1189 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1190 (*II).getOpcode() == TargetOpcode::COPY ||
1191 (*II).getOpcode() == SPIRV::OpVariable) &&
1192 getImm(
I.getOperand(2), MRI));
1194 bool IsGVInit =
false;
1198 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1199 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1200 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1201 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1211 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1224 "incompatible result and operand types in a bitcast");
1226 MachineInstrBuilder MIB =
1227 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1234 : SPIRV::OpInBoundsPtrAccessChain))
1238 .
addUse(
I.getOperand(2).getReg())
1241 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1245 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1247 .
addUse(
I.getOperand(2).getReg())
1256 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1259 .
addImm(
static_cast<uint32_t
>(
1260 SPIRV::Opcode::InBoundsPtrAccessChain))
1263 .
addUse(
I.getOperand(2).getReg());
1268 case TargetOpcode::G_ATOMICRMW_OR:
1269 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1270 case TargetOpcode::G_ATOMICRMW_ADD:
1271 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1272 case TargetOpcode::G_ATOMICRMW_AND:
1273 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1274 case TargetOpcode::G_ATOMICRMW_MAX:
1275 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1276 case TargetOpcode::G_ATOMICRMW_MIN:
1277 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1278 case TargetOpcode::G_ATOMICRMW_SUB:
1279 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1280 case TargetOpcode::G_ATOMICRMW_XOR:
1281 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1282 case TargetOpcode::G_ATOMICRMW_UMAX:
1283 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1284 case TargetOpcode::G_ATOMICRMW_UMIN:
1285 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1286 case TargetOpcode::G_ATOMICRMW_XCHG:
1287 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1288 case TargetOpcode::G_ATOMIC_CMPXCHG:
1289 return selectAtomicCmpXchg(ResVReg, ResType,
I);
1291 case TargetOpcode::G_ATOMICRMW_FADD:
1292 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1293 case TargetOpcode::G_ATOMICRMW_FSUB:
1295 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1296 ResType->
getOpcode() == SPIRV::OpTypeVector
1298 : SPIRV::OpFNegate);
1299 case TargetOpcode::G_ATOMICRMW_FMIN:
1300 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1301 case TargetOpcode::G_ATOMICRMW_FMAX:
1302 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1304 case TargetOpcode::G_FENCE:
1305 return selectFence(
I);
1307 case TargetOpcode::G_STACKSAVE:
1308 return selectStackSave(ResVReg, ResType,
I);
1309 case TargetOpcode::G_STACKRESTORE:
1310 return selectStackRestore(
I);
1312 case TargetOpcode::G_UNMERGE_VALUES:
1318 case TargetOpcode::G_TRAP:
1319 case TargetOpcode::G_UBSANTRAP:
1320 case TargetOpcode::DBG_LABEL:
1322 case TargetOpcode::G_DEBUGTRAP:
1323 return selectDebugTrap(ResVReg, ResType,
I);
1330bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1331 SPIRVTypeInst ResType,
1332 MachineInstr &
I)
const {
1333 unsigned Opcode = SPIRV::OpNop;
1340bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1341 SPIRVTypeInst ResType,
1343 GL::GLSLExtInst GLInst,
1344 bool setMIFlags,
bool useMISrc,
1347 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1348 std::string DiagMsg;
1349 raw_string_ostream OS(DiagMsg);
1350 I.print(OS,
true,
false,
false,
false);
1351 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1354 return selectExtInst(ResVReg, ResType,
I,
1355 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1356 setMIFlags, useMISrc, SrcRegs);
1359bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1360 SPIRVTypeInst ResType,
1362 CL::OpenCLExtInst CLInst,
1363 bool setMIFlags,
bool useMISrc,
1365 return selectExtInst(ResVReg, ResType,
I,
1366 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1367 setMIFlags, useMISrc, SrcRegs);
1370bool SPIRVInstructionSelector::selectExtInst(
1371 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1372 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1374 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1375 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1376 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1380bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1381 SPIRVTypeInst ResType,
1384 bool setMIFlags,
bool useMISrc,
1387 for (
const auto &[InstructionSet, Opcode] : Insts) {
1391 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1394 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1399 const unsigned NumOps =
I.getNumOperands();
1402 I.getOperand(Index).getType() ==
1403 MachineOperand::MachineOperandType::MO_IntrinsicID)
1406 MIB.
add(
I.getOperand(Index));
1418bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1419 SPIRVTypeInst ResType,
1420 MachineInstr &
I)
const {
1421 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1422 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1423 for (
const auto &Ex : ExtInsts) {
1424 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1425 uint32_t Opcode = Ex.second;
1429 MachineIRBuilder MIRBuilder(
I);
1432 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1437 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1440 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1443 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1446 .
addImm(
static_cast<uint32_t
>(Ex.first))
1448 .
add(
I.getOperand(2))
1452 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1453 .
addDef(
I.getOperand(1).getReg())
1462bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1463 SPIRVTypeInst ResType,
1464 MachineInstr &
I)
const {
1465 Register CosResVReg =
I.getOperand(1).getReg();
1466 unsigned SrcIdx =
I.getNumExplicitDefs();
1471 MachineIRBuilder MIRBuilder(
I);
1473 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1478 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1481 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1483 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1486 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1488 .
add(
I.getOperand(SrcIdx))
1491 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1499 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1502 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1504 .
add(
I.getOperand(SrcIdx))
1506 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1509 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1511 .
add(
I.getOperand(SrcIdx))
1518bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1519 SPIRVTypeInst ResType,
1521 std::vector<Register> Srcs,
1522 unsigned Opcode)
const {
1523 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1533bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1534 SPIRVTypeInst ResType,
1536 unsigned Opcode)
const {
1538 Register SrcReg =
I.getOperand(1).getReg();
1543 unsigned DefOpCode = DefIt->getOpcode();
1544 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1547 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1548 DefOpCode = VRD->getOpcode();
1550 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1551 DefOpCode == TargetOpcode::G_CONSTANT ||
1552 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1558 uint32_t SpecOpcode = 0;
1560 case SPIRV::OpConvertPtrToU:
1561 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1563 case SPIRV::OpConvertUToPtr:
1564 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1569 TII.get(SPIRV::OpSpecConstantOp))
1579 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1583bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1584 SPIRVTypeInst ResType,
1585 MachineInstr &
I)
const {
1586 Register OpReg =
I.getOperand(1).getReg();
1587 SPIRVTypeInst OpType =
1591 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1601 if (
MemOp->isVolatile())
1602 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1603 if (
MemOp->isNonTemporal())
1604 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1606 if (!ST->isShader() &&
MemOp->getAlign().value())
1607 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1611 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1612 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1616 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1618 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1622 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1626 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1628 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1640 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1642 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1644 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1648bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1649 SPIRVTypeInst ResType,
1650 MachineInstr &
I)
const {
1652 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1657 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1658 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1660 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1664 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1668 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1669 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1670 I.getDebugLoc(),
I);
1674 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1678 if (!
I.getNumMemOperands()) {
1679 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1681 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1684 MachineIRBuilder MIRBuilder(
I);
1691bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1693 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1694 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1699 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1700 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1705 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1709 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1710 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1711 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1712 TII.get(SPIRV::OpImageWrite))
1718 if (sampledTypeIsSignedInteger(LLVMHandleType))
1721 BMI.constrainAllUses(
TII,
TRI, RBI);
1727 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpStore))
1730 if (!
I.getNumMemOperands()) {
1731 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1733 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1736 MachineIRBuilder MIRBuilder(
I);
1743bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
1744 SPIRVTypeInst ResType,
1745 MachineInstr &
I)
const {
1746 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
1754 Register PtrsReg =
I.getOperand(2).getReg();
1755 uint32_t Alignment =
I.getOperand(3).getImm();
1756 Register MaskReg =
I.getOperand(4).getReg();
1757 Register PassthruReg =
I.getOperand(5).getReg();
1758 Register AlignmentReg = buildI32Constant(Alignment,
I);
1762 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
1773bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
1774 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
1781 Register ValuesReg =
I.getOperand(1).getReg();
1782 Register PtrsReg =
I.getOperand(2).getReg();
1783 uint32_t Alignment =
I.getOperand(3).getImm();
1784 Register MaskReg =
I.getOperand(4).getReg();
1785 Register AlignmentReg = buildI32Constant(Alignment,
I);
1789 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
1798bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
1799 const Twine &Msg)
const {
1800 const Function &
F =
I.getMF()->getFunction();
1801 F.getContext().diagnose(
1802 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
1806bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
1807 SPIRVTypeInst ResType,
1808 MachineInstr &
I)
const {
1809 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1811 "llvm.stacksave intrinsic: this instruction requires the following "
1812 "SPIR-V extension: SPV_INTEL_variable_length_array",
1815 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
1822bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
1823 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1825 "llvm.stackrestore intrinsic: this instruction requires the following "
1826 "SPIR-V extension: SPV_INTEL_variable_length_array",
1828 if (!
I.getOperand(0).isReg())
1831 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
1832 .
addUse(
I.getOperand(0).getReg())
1838SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
1839 MachineIRBuilder MIRBuilder(
I);
1840 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1847 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
1851 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1852 Type *ArrTy = ArrayType::get(ValTy, Num);
1854 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
1857 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
1864 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
1867 .
addImm(SPIRV::StorageClass::UniformConstant)
1878bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
1881 Register DstReg =
I.getOperand(0).getReg();
1886 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
1891 "Unable to determine pointee type size for OpCopyMemory");
1892 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
1893 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
1895 "OpCopyMemory requires the size to match the pointee type size");
1896 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
1899 if (
I.getNumMemOperands()) {
1900 MachineIRBuilder MIRBuilder(
I);
1907bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
1910 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
1911 .
addUse(
I.getOperand(0).getReg())
1913 .
addUse(
I.getOperand(2).getReg());
1914 if (
I.getNumMemOperands()) {
1915 MachineIRBuilder MIRBuilder(
I);
1922bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
1923 MachineInstr &
I)
const {
1924 Register SrcReg =
I.getOperand(1).getReg();
1925 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
1926 Register VarReg = getOrCreateMemSetGlobal(
I);
1929 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1931 ValTy,
I, SPIRV::StorageClass::UniformConstant);
1933 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
1937 if (!selectCopyMemory(
I, SrcReg))
1940 if (!selectCopyMemorySized(
I, SrcReg))
1943 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
1944 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
1949bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
1950 SPIRVTypeInst ResType,
1953 unsigned NegateOpcode)
const {
1955 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1958 Register ScopeReg = buildI32Constant(Scope,
I);
1960 Register Ptr =
I.getOperand(1).getReg();
1966 Register MemSemReg = buildI32Constant(MemSem ,
I);
1968 Register ValueReg =
I.getOperand(2).getReg();
1969 if (NegateOpcode != 0) {
1972 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
1977 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
1988bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
1989 unsigned ArgI =
I.getNumOperands() - 1;
1991 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
1992 SPIRVTypeInst SrcType =
1994 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
1996 "cannot select G_UNMERGE_VALUES with a non-vector argument");
1998 SPIRVTypeInst ScalarType =
2001 unsigned CurrentIndex = 0;
2002 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2003 Register ResVReg =
I.getOperand(i).getReg();
2006 LLT ResLLT = MRI->
getType(ResVReg);
2012 ResType = ScalarType;
2018 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2021 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2027 for (
unsigned j = 0;
j < NumElements; ++
j) {
2028 MIB.
addImm(CurrentIndex + j);
2030 CurrentIndex += NumElements;
2034 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2046bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2049 Register MemSemReg = buildI32Constant(MemSem,
I);
2051 uint32_t
Scope =
static_cast<uint32_t
>(
2053 Register ScopeReg = buildI32Constant(Scope,
I);
2055 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2062bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2063 SPIRVTypeInst ResType,
2065 unsigned Opcode)
const {
2066 Type *ResTy =
nullptr;
2070 "Not enough info to select the arithmetic with overflow instruction");
2073 "with overflow instruction");
2079 MachineIRBuilder MIRBuilder(
I);
2081 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2082 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2088 Register ZeroReg = buildZerosVal(ResType,
I);
2093 if (ResName.
size() > 0)
2098 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2101 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2102 MIB.
addUse(
I.getOperand(i).getReg());
2107 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2108 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2110 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2111 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2118 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2119 .
addDef(
I.getOperand(1).getReg())
2127bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2128 SPIRVTypeInst ResType,
2129 MachineInstr &
I)
const {
2133 Register Ptr =
I.getOperand(2).getReg();
2136 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2139 ScopeReg = buildI32Constant(Scope,
I);
2141 unsigned ScSem =
static_cast<uint32_t
>(
2144 unsigned MemSemEq =
static_cast<uint32_t
>(
getMemSemantics(AO)) | ScSem;
2145 Register MemSemEqReg = buildI32Constant(MemSemEq,
I);
2147 unsigned MemSemNeq =
static_cast<uint32_t
>(
getMemSemantics(FO)) | ScSem;
2148 if (MemSemEq == MemSemNeq)
2149 MemSemNeqReg = MemSemEqReg;
2151 MemSemNeqReg = buildI32Constant(MemSemEq,
I);
2154 ScopeReg =
I.getOperand(5).getReg();
2155 MemSemEqReg =
I.getOperand(6).getReg();
2156 MemSemNeqReg =
I.getOperand(7).getReg();
2160 Register Val =
I.getOperand(4).getReg();
2164 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2183 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2190 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2202 case SPIRV::StorageClass::DeviceOnlyINTEL:
2203 case SPIRV::StorageClass::HostOnlyINTEL:
2212 bool IsGRef =
false;
2213 bool IsAllowedRefs =
2215 unsigned Opcode = It.getOpcode();
2216 if (Opcode == SPIRV::OpConstantComposite ||
2217 Opcode == SPIRV::OpVariable ||
2218 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2219 return IsGRef = true;
2220 return Opcode == SPIRV::OpName;
2222 return IsAllowedRefs && IsGRef;
2225Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2226 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2228 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2232SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2234 uint32_t Opcode)
const {
2235 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2236 TII.get(SPIRV::OpSpecConstantOp))
2244SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2245 SPIRVTypeInst SrcPtrTy)
const {
2246 SPIRVTypeInst GenericPtrTy =
2250 SPIRV::StorageClass::Generic),
2252 MachineFunction *MF =
I.getParent()->getParent();
2254 MachineInstrBuilder MIB = buildSpecConstantOp(
2256 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2266bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2267 SPIRVTypeInst ResType,
2268 MachineInstr &
I)
const {
2272 Register SrcPtr =
I.getOperand(1).getReg();
2276 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2277 ResType->
getOpcode() != SPIRV::OpTypePointer)
2278 return BuildCOPY(ResVReg, SrcPtr,
I);
2288 unsigned SpecOpcode =
2290 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2293 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2300 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2302 .constrainAllUses(
TII,
TRI, RBI);
2304 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2306 buildSpecConstantOp(
2308 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2309 .constrainAllUses(
TII,
TRI, RBI);
2316 return BuildCOPY(ResVReg, SrcPtr,
I);
2318 if ((SrcSC == SPIRV::StorageClass::Function &&
2319 DstSC == SPIRV::StorageClass::Private) ||
2320 (DstSC == SPIRV::StorageClass::Function &&
2321 SrcSC == SPIRV::StorageClass::Private))
2322 return BuildCOPY(ResVReg, SrcPtr,
I);
2326 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2329 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2332 SPIRVTypeInst GenericPtrTy =
2351 return selectUnOp(ResVReg, ResType,
I,
2352 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2354 return selectUnOp(ResVReg, ResType,
I,
2355 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2357 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2359 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2369 return SPIRV::OpFOrdEqual;
2371 return SPIRV::OpFOrdGreaterThanEqual;
2373 return SPIRV::OpFOrdGreaterThan;
2375 return SPIRV::OpFOrdLessThanEqual;
2377 return SPIRV::OpFOrdLessThan;
2379 return SPIRV::OpFOrdNotEqual;
2381 return SPIRV::OpOrdered;
2383 return SPIRV::OpFUnordEqual;
2385 return SPIRV::OpFUnordGreaterThanEqual;
2387 return SPIRV::OpFUnordGreaterThan;
2389 return SPIRV::OpFUnordLessThanEqual;
2391 return SPIRV::OpFUnordLessThan;
2393 return SPIRV::OpFUnordNotEqual;
2395 return SPIRV::OpUnordered;
2405 return SPIRV::OpIEqual;
2407 return SPIRV::OpINotEqual;
2409 return SPIRV::OpSGreaterThanEqual;
2411 return SPIRV::OpSGreaterThan;
2413 return SPIRV::OpSLessThanEqual;
2415 return SPIRV::OpSLessThan;
2417 return SPIRV::OpUGreaterThanEqual;
2419 return SPIRV::OpUGreaterThan;
2421 return SPIRV::OpULessThanEqual;
2423 return SPIRV::OpULessThan;
2432 return SPIRV::OpPtrEqual;
2434 return SPIRV::OpPtrNotEqual;
2445 return SPIRV::OpLogicalEqual;
2447 return SPIRV::OpLogicalNotEqual;
2481bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2482 SPIRVTypeInst ResType,
2484 unsigned OpAnyOrAll)
const {
2485 assert(
I.getNumOperands() == 3);
2486 assert(
I.getOperand(2).isReg());
2488 Register InputRegister =
I.getOperand(2).getReg();
2491 assert(InputType &&
"VReg has no type assigned");
2494 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2495 if (IsBoolTy && !IsVectorTy) {
2496 assert(ResVReg ==
I.getOperand(0).getReg());
2497 return BuildCOPY(ResVReg, InputRegister,
I);
2501 unsigned SpirvNotEqualId =
2502 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2504 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2509 IsBoolTy ? InputRegister
2517 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2519 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2536bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2537 SPIRVTypeInst ResType,
2538 MachineInstr &
I)
const {
2539 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2542bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2543 SPIRVTypeInst ResType,
2544 MachineInstr &
I)
const {
2545 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2549bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2550 SPIRVTypeInst ResType,
2551 MachineInstr &
I)
const {
2552 assert(
I.getNumOperands() == 4);
2553 assert(
I.getOperand(2).isReg());
2554 assert(
I.getOperand(3).isReg());
2556 [[maybe_unused]] SPIRVTypeInst VecType =
2561 "dot product requires a vector of at least 2 components");
2563 [[maybe_unused]] SPIRVTypeInst EltType =
2572 .
addUse(
I.getOperand(2).getReg())
2573 .
addUse(
I.getOperand(3).getReg())
2578bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2579 SPIRVTypeInst ResType,
2582 assert(
I.getNumOperands() == 4);
2583 assert(
I.getOperand(2).isReg());
2584 assert(
I.getOperand(3).isReg());
2587 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2591 .
addUse(
I.getOperand(2).getReg())
2592 .
addUse(
I.getOperand(3).getReg())
2599bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2600 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2601 assert(
I.getNumOperands() == 4);
2602 assert(
I.getOperand(2).isReg());
2603 assert(
I.getOperand(3).isReg());
2607 Register Vec0 =
I.getOperand(2).getReg();
2608 Register Vec1 =
I.getOperand(3).getReg();
2612 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2621 "dot product requires a vector of at least 2 components");
2624 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2634 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2645 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2657bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2658 SPIRVTypeInst ResType,
2659 MachineInstr &
I)
const {
2661 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2664 .
addUse(
I.getOperand(2).getReg())
2669bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2670 SPIRVTypeInst ResType,
2671 MachineInstr &
I)
const {
2673 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2676 .
addUse(
I.getOperand(2).getReg())
2681template <
bool Signed>
2682bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2683 SPIRVTypeInst ResType,
2684 MachineInstr &
I)
const {
2685 assert(
I.getNumOperands() == 5);
2686 assert(
I.getOperand(2).isReg());
2687 assert(
I.getOperand(3).isReg());
2688 assert(
I.getOperand(4).isReg());
2691 Register Acc =
I.getOperand(2).getReg();
2695 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2697 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
2702 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
2705 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2717template <
bool Signed>
2718bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
2719 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2720 assert(
I.getNumOperands() == 5);
2721 assert(
I.getOperand(2).isReg());
2722 assert(
I.getOperand(3).isReg());
2723 assert(
I.getOperand(4).isReg());
2726 Register Acc =
I.getOperand(2).getReg();
2732 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
2736 for (
unsigned i = 0; i < 4; i++) {
2759 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
2779 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2794bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
2795 SPIRVTypeInst ResType,
2796 MachineInstr &
I)
const {
2797 assert(
I.getNumOperands() == 3);
2798 assert(
I.getOperand(2).isReg());
2800 Register VZero = buildZerosValF(ResType,
I);
2801 Register VOne = buildOnesValF(ResType,
I);
2803 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
2806 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2808 .
addUse(
I.getOperand(2).getReg())
2815bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
2816 SPIRVTypeInst ResType,
2817 MachineInstr &
I)
const {
2818 assert(
I.getNumOperands() == 3);
2819 assert(
I.getOperand(2).isReg());
2821 Register InputRegister =
I.getOperand(2).getReg();
2823 auto &
DL =
I.getDebugLoc();
2833 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
2835 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
2843 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2848 if (NeedsConversion) {
2849 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
2860bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
2861 SPIRVTypeInst ResType,
2863 unsigned Opcode)
const {
2867 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2873 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
2874 BMI.addUse(
I.getOperand(J).getReg());
2881bool SPIRVInstructionSelector::selectWaveActiveCountBits(
2882 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2887 if (!selectWaveOpInst(BallotReg, BallotType,
I,
2888 SPIRV::OpGroupNonUniformBallot))
2893 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2898 .
addImm(SPIRV::GroupOperation::Reduce)
2907 if (
Type->getOpcode() != SPIRV::OpTypeVector)
2911 return Type->getOperand(2).getImm();
2914bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
2915 SPIRVTypeInst ResType,
2916 MachineInstr &
I)
const {
2921 Register InputReg =
I.getOperand(2).getReg();
2926 bool IsVector = NumElems > 1;
2929 SPIRVTypeInst ElemInputType = InputType;
2930 SPIRVTypeInst ElemBoolType = ResType;
2943 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
2944 SPIRV::OpGroupNonUniformAllEqual);
2949 ElementResults.
reserve(NumElems);
2951 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
2964 ElemInput = Extracted;
2970 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
2981 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
2992bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
2993 SPIRVTypeInst ResType,
2994 MachineInstr &
I)
const {
2996 assert(
I.getNumOperands() == 3);
2998 auto Op =
I.getOperand(2);
3010 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3032 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3036 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3043bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3044 SPIRVTypeInst ResType,
3046 bool IsUnsigned)
const {
3047 return selectWaveReduce(
3048 ResVReg, ResType,
I, IsUnsigned,
3049 [&](
Register InputRegister,
bool IsUnsigned) {
3050 const bool IsFloatTy =
3052 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3053 : SPIRV::OpGroupNonUniformSMax;
3054 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3058bool SPIRVInstructionSelector::selectWaveReduceMin(
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::OpGroupNonUniformUMin
3068 : SPIRV::OpGroupNonUniformSMin;
3069 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3073bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3074 SPIRVTypeInst ResType,
3075 MachineInstr &
I)
const {
3076 return selectWaveReduce(ResVReg, ResType,
I,
false,
3077 [&](
Register InputRegister,
bool IsUnsigned) {
3079 InputRegister, SPIRV::OpTypeFloat);
3080 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3081 : SPIRV::OpGroupNonUniformIAdd;
3085bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3086 SPIRVTypeInst ResType,
3087 MachineInstr &
I)
const {
3088 return selectWaveReduce(ResVReg, ResType,
I,
false,
3089 [&](
Register InputRegister,
bool IsUnsigned) {
3091 InputRegister, SPIRV::OpTypeFloat);
3092 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3093 : SPIRV::OpGroupNonUniformIMul;
3097template <
typename PickOpcodeFn>
3098bool SPIRVInstructionSelector::selectWaveReduce(
3099 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3100 PickOpcodeFn &&PickOpcode)
const {
3101 assert(
I.getNumOperands() == 3);
3102 assert(
I.getOperand(2).isReg());
3104 Register InputRegister =
I.getOperand(2).getReg();
3111 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3117 .
addImm(SPIRV::GroupOperation::Reduce)
3118 .
addUse(
I.getOperand(2).getReg())
3123bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3124 SPIRVTypeInst ResType,
3126 unsigned Opcode)
const {
3127 return selectWaveReduce(
3128 ResVReg, ResType,
I,
false,
3129 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3132bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3133 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3134 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3135 [&](
Register InputRegister,
bool IsUnsigned) {
3137 InputRegister, SPIRV::OpTypeFloat);
3139 ? SPIRV::OpGroupNonUniformFAdd
3140 : SPIRV::OpGroupNonUniformIAdd;
3144bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3145 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3146 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3147 [&](
Register InputRegister,
bool IsUnsigned) {
3149 InputRegister, SPIRV::OpTypeFloat);
3151 ? SPIRV::OpGroupNonUniformFMul
3152 : SPIRV::OpGroupNonUniformIMul;
3156template <
typename PickOpcodeFn>
3157bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3158 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3159 PickOpcodeFn &&PickOpcode)
const {
3160 assert(
I.getNumOperands() == 3);
3161 assert(
I.getOperand(2).isReg());
3163 Register InputRegister =
I.getOperand(2).getReg();
3170 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3176 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3177 .
addUse(
I.getOperand(2).getReg())
3182bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3183 SPIRVTypeInst ResType,
3184 MachineInstr &
I)
const {
3186 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3189 .
addUse(
I.getOperand(1).getReg())
3194bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3195 SPIRVTypeInst ResType,
3196 MachineInstr &
I)
const {
3202 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3204 Register OpReg =
I.getOperand(1).getReg();
3205 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
3206 if (
Def->getOpcode() == TargetOpcode::COPY)
3209 switch (
Def->getOpcode()) {
3210 case SPIRV::ASSIGN_TYPE:
3211 if (MachineInstr *AssignToDef =
3213 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3214 Reg =
Def->getOperand(2).getReg();
3217 case SPIRV::OpUndef:
3218 Reg =
Def->getOperand(1).getReg();
3221 unsigned DestOpCode;
3223 DestOpCode = SPIRV::OpConstantNull;
3225 DestOpCode = TargetOpcode::COPY;
3228 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3229 .
addDef(
I.getOperand(0).getReg())
3237bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3238 SPIRVTypeInst ResType,
3239 MachineInstr &
I)
const {
3241 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3243 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3247 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3252 for (
unsigned i =
I.getNumExplicitDefs();
3253 i <
I.getNumExplicitOperands() && IsConst; ++i)
3257 if (!IsConst &&
N < 2)
3259 "There must be at least two constituent operands in a vector");
3262 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3263 TII.get(IsConst ? SPIRV::OpConstantComposite
3264 : SPIRV::OpCompositeConstruct))
3267 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3268 MIB.
addUse(
I.getOperand(i).getReg());
3273bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3274 SPIRVTypeInst ResType,
3275 MachineInstr &
I)
const {
3277 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3279 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3285 if (!
I.getOperand(
OpIdx).isReg())
3292 if (!IsConst &&
N < 2)
3294 "There must be at least two constituent operands in a vector");
3297 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3298 TII.get(IsConst ? SPIRV::OpConstantComposite
3299 : SPIRV::OpCompositeConstruct))
3302 for (
unsigned i = 0; i <
N; ++i)
3308bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3309 SPIRVTypeInst ResType,
3310 MachineInstr &
I)
const {
3315 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3317 Opcode = SPIRV::OpDemoteToHelperInvocation;
3319 Opcode = SPIRV::OpKill;
3321 if (MachineInstr *NextI =
I.getNextNode()) {
3323 NextI->eraseFromParent();
3333bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3334 SPIRVTypeInst ResType,
unsigned CmpOpc,
3335 MachineInstr &
I)
const {
3336 Register Cmp0 =
I.getOperand(2).getReg();
3337 Register Cmp1 =
I.getOperand(3).getReg();
3340 "CMP operands should have the same type");
3341 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3351bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3352 SPIRVTypeInst ResType,
3353 MachineInstr &
I)
const {
3354 auto Pred =
I.getOperand(1).getPredicate();
3357 Register CmpOperand =
I.getOperand(2).getReg();
3364 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3368SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3369 SPIRVTypeInst ResType)
const {
3371 SPIRVTypeInst SpvI32Ty =
3374 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3381 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3384 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3387 .
addImm(APInt(32, Val).getZExtValue());
3389 GR.
add(ConstInt,
MI);
3394bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
3395 SPIRVTypeInst ResType,
3396 MachineInstr &
I)
const {
3398 return selectCmp(ResVReg, ResType, CmpOp,
I);
3401bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
3402 SPIRVTypeInst ResType,
3403 MachineInstr &
I)
const {
3405 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
3412 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
3413 ResType->
getOpcode() != SPIRV::OpTypeFloat)
3416 MachineIRBuilder MIRBuilder(
I);
3418 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
3424 "only float operands supported by GLSL extended math");
3427 MIRBuilder, SpirvScalarType);
3429 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
3430 ? SPIRV::OpVectorTimesScalar
3433 if (!selectOpWithSrcs(ArgReg, ResType,
I,
3434 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
3436 if (!selectExtInst(ResVReg, ResType,
I,
3437 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
3447Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
3448 MachineInstr &
I)
const {
3451 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3456bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
3462 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
3470 if (
Def->getOpcode() == SPIRV::OpConstantNull)
3473 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
3474 Def->getOpcode() == SPIRV::OpConstantI)
3487 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
3488 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
3490 Intrinsic::spv_const_composite)) {
3491 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
3492 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
3493 if (!IsZero(
Def->getOperand(i).getReg()))
3502Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
3503 MachineInstr &
I)
const {
3507 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3512Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
3513 MachineInstr &
I)
const {
3517 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3523 SPIRVTypeInst ResType,
3524 MachineInstr &
I)
const {
3528 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3533bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
3534 SPIRVTypeInst ResType,
3535 MachineInstr &
I)
const {
3536 Register SelectFirstArg =
I.getOperand(2).getReg();
3537 Register SelectSecondArg =
I.getOperand(3).getReg();
3546 SPIRV::OpTypeVector;
3553 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
3554 }
else if (IsPtrTy) {
3555 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
3557 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
3561 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
3562 }
else if (IsPtrTy) {
3563 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
3565 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3568 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3571 .
addUse(
I.getOperand(1).getReg())
3580bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
3581 SPIRVTypeInst ResType,
3583 MachineInstr &InsertAt,
3584 bool IsSigned)
const {
3586 Register ZeroReg = buildZerosVal(ResType, InsertAt);
3587 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
3588 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
3590 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3602bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
3603 SPIRVTypeInst ResType,
3604 MachineInstr &
I,
bool IsSigned,
3605 unsigned Opcode)
const {
3606 Register SrcReg =
I.getOperand(1).getReg();
3612 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
3617 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
3619 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
3622bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
3623 SPIRVTypeInst ResType, MachineInstr &
I,
3624 bool IsSigned)
const {
3625 Register SrcReg =
I.getOperand(1).getReg();
3627 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
3631 if (ResType == SrcType)
3632 return BuildCOPY(ResVReg, SrcReg,
I);
3634 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3635 return selectUnOp(ResVReg, ResType,
I, Opcode);
3638bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
3639 SPIRVTypeInst ResType,
3641 bool IsSigned)
const {
3642 MachineIRBuilder MIRBuilder(
I);
3643 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
3658 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
3661 .
addUse(
I.getOperand(1).getReg())
3662 .
addUse(
I.getOperand(2).getReg())
3668 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
3671 .
addUse(
I.getOperand(1).getReg())
3672 .
addUse(
I.getOperand(2).getReg())
3680 unsigned SelectOpcode =
3681 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
3686 .
addUse(buildOnesVal(
true, ResType,
I))
3687 .
addUse(buildZerosVal(ResType,
I))
3694 .
addUse(buildOnesVal(
false, ResType,
I))
3699bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
3702 SPIRVTypeInst IntTy,
3703 SPIRVTypeInst BoolTy)
const {
3706 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
3707 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
3709 Register One = buildOnesVal(
false, IntTy,
I);
3717 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
3726bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
3727 SPIRVTypeInst ResType,
3728 MachineInstr &
I)
const {
3729 Register IntReg =
I.getOperand(1).getReg();
3732 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
3733 if (ArgType == ResType)
3734 return BuildCOPY(ResVReg, IntReg,
I);
3736 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3737 return selectUnOp(ResVReg, ResType,
I, Opcode);
3740bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
3741 SPIRVTypeInst ResType,
3742 MachineInstr &
I)
const {
3743 unsigned Opcode =
I.getOpcode();
3744 unsigned TpOpcode = ResType->
getOpcode();
3746 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
3747 assert(Opcode == TargetOpcode::G_CONSTANT &&
3748 I.getOperand(1).getCImm()->isZero());
3749 MachineBasicBlock &DepMBB =
I.getMF()->front();
3752 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
3759 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
3762bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
3763 SPIRVTypeInst ResType,
3764 MachineInstr &
I)
const {
3765 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3772bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
3773 SPIRVTypeInst ResType,
3774 MachineInstr &
I)
const {
3776 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
3780 .
addUse(
I.getOperand(3).getReg())
3782 .
addUse(
I.getOperand(2).getReg());
3783 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
3789bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
3790 SPIRVTypeInst ResType,
3791 MachineInstr &
I)
const {
3792 Type *MaybeResTy =
nullptr;
3797 "Expected aggregate type for extractv instruction");
3799 SPIRV::AccessQualifier::ReadWrite,
false);
3803 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3806 .
addUse(
I.getOperand(2).getReg());
3807 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
3813bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
3814 SPIRVTypeInst ResType,
3815 MachineInstr &
I)
const {
3816 if (
getImm(
I.getOperand(4), MRI))
3817 return selectInsertVal(ResVReg, ResType,
I);
3819 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
3822 .
addUse(
I.getOperand(2).getReg())
3823 .
addUse(
I.getOperand(3).getReg())
3824 .
addUse(
I.getOperand(4).getReg())
3829bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
3830 SPIRVTypeInst ResType,
3831 MachineInstr &
I)
const {
3832 if (
getImm(
I.getOperand(3), MRI))
3833 return selectExtractVal(ResVReg, ResType,
I);
3835 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
3838 .
addUse(
I.getOperand(2).getReg())
3839 .
addUse(
I.getOperand(3).getReg())
3844bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
3845 SPIRVTypeInst ResType,
3846 MachineInstr &
I)
const {
3847 const bool IsGEPInBounds =
I.getOperand(2).getImm();
3853 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
3854 : SPIRV::OpAccessChain)
3855 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
3856 :
SPIRV::OpPtrAccessChain);
3858 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3862 .
addUse(
I.getOperand(3).getReg());
3864 (Opcode == SPIRV::OpPtrAccessChain ||
3865 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
3866 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
3867 "Cannot translate GEP to OpAccessChain. First index must be 0.");
3870 const unsigned StartingIndex =
3871 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
3874 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
3875 Res.addUse(
I.getOperand(i).getReg());
3876 Res.constrainAllUses(
TII,
TRI, RBI);
3881bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
3883 unsigned Lim =
I.getNumExplicitOperands();
3884 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
3885 Register OpReg =
I.getOperand(i).getReg();
3886 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
3888 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
3889 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
3890 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
3897 MachineFunction *MF =
I.getMF();
3909 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3910 TII.get(SPIRV::OpSpecConstantOp))
3913 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
3915 GR.
add(OpDefine, MIB);
3921bool SPIRVInstructionSelector::selectDerivativeInst(
3922 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
3923 const unsigned DPdOpCode)
const {
3926 errorIfInstrOutsideShader(
I);
3931 Register SrcReg =
I.getOperand(2).getReg();
3936 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
3939 .
addUse(
I.getOperand(2).getReg());
3941 MachineIRBuilder MIRBuilder(
I);
3944 if (componentCount != 1)
3948 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
3952 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
3957 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
3962 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
3970bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
3971 SPIRVTypeInst ResType,
3972 MachineInstr &
I)
const {
3976 case Intrinsic::spv_load:
3977 return selectLoad(ResVReg, ResType,
I);
3978 case Intrinsic::spv_store:
3979 return selectStore(
I);
3980 case Intrinsic::spv_extractv:
3981 return selectExtractVal(ResVReg, ResType,
I);
3982 case Intrinsic::spv_insertv:
3983 return selectInsertVal(ResVReg, ResType,
I);
3984 case Intrinsic::spv_extractelt:
3985 return selectExtractElt(ResVReg, ResType,
I);
3986 case Intrinsic::spv_insertelt:
3987 return selectInsertElt(ResVReg, ResType,
I);
3988 case Intrinsic::spv_gep:
3989 return selectGEP(ResVReg, ResType,
I);
3990 case Intrinsic::spv_bitcast: {
3991 Register OpReg =
I.getOperand(2).getReg();
3992 SPIRVTypeInst OpType =
3996 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
3998 case Intrinsic::spv_unref_global:
3999 case Intrinsic::spv_init_global: {
4000 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4005 Register GVarVReg =
MI->getOperand(0).getReg();
4006 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4011 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4013 MI->eraseFromParent();
4017 case Intrinsic::spv_undef: {
4018 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4024 case Intrinsic::spv_const_composite: {
4026 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4032 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4034 MachineIRBuilder MIR(
I);
4036 MIR, SPIRV::OpConstantComposite, 3,
4037 SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,
4039 for (
auto *Instr : Instructions) {
4040 Instr->setDebugLoc(
I.getDebugLoc());
4045 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4052 case Intrinsic::spv_assign_name: {
4053 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4054 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4055 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4056 i <
I.getNumExplicitOperands(); ++i) {
4057 MIB.
addImm(
I.getOperand(i).getImm());
4062 case Intrinsic::spv_switch: {
4063 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4064 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4065 if (
I.getOperand(i).isReg())
4066 MIB.
addReg(
I.getOperand(i).getReg());
4067 else if (
I.getOperand(i).isCImm())
4068 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4069 else if (
I.getOperand(i).isMBB())
4070 MIB.
addMBB(
I.getOperand(i).getMBB());
4077 case Intrinsic::spv_loop_merge: {
4078 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4079 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4080 if (
I.getOperand(i).isMBB())
4081 MIB.
addMBB(
I.getOperand(i).getMBB());
4088 case Intrinsic::spv_loop_control_intel: {
4090 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4091 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4096 case Intrinsic::spv_selection_merge: {
4098 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4099 assert(
I.getOperand(1).isMBB() &&
4100 "operand 1 to spv_selection_merge must be a basic block");
4101 MIB.
addMBB(
I.getOperand(1).getMBB());
4102 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4106 case Intrinsic::spv_cmpxchg:
4107 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4108 case Intrinsic::spv_unreachable:
4109 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4112 case Intrinsic::spv_alloca:
4113 return selectFrameIndex(ResVReg, ResType,
I);
4114 case Intrinsic::spv_alloca_array:
4115 return selectAllocaArray(ResVReg, ResType,
I);
4116 case Intrinsic::spv_assume:
4118 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4119 .
addUse(
I.getOperand(1).getReg())
4124 case Intrinsic::spv_expect:
4126 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4129 .
addUse(
I.getOperand(2).getReg())
4130 .
addUse(
I.getOperand(3).getReg())
4135 case Intrinsic::arithmetic_fence:
4136 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4137 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4140 .
addUse(
I.getOperand(2).getReg())
4144 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4146 case Intrinsic::spv_thread_id:
4152 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4154 case Intrinsic::spv_thread_id_in_group:
4160 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4162 case Intrinsic::spv_group_id:
4168 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4170 case Intrinsic::spv_flattened_thread_id_in_group:
4177 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4179 case Intrinsic::spv_workgroup_size:
4180 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4182 case Intrinsic::spv_global_size:
4183 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4185 case Intrinsic::spv_global_offset:
4186 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4188 case Intrinsic::spv_num_workgroups:
4189 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4191 case Intrinsic::spv_subgroup_size:
4192 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4194 case Intrinsic::spv_num_subgroups:
4195 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4197 case Intrinsic::spv_subgroup_id:
4198 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4199 case Intrinsic::spv_subgroup_local_invocation_id:
4200 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4201 ResVReg, ResType,
I);
4202 case Intrinsic::spv_subgroup_max_size:
4203 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4205 case Intrinsic::spv_fdot:
4206 return selectFloatDot(ResVReg, ResType,
I);
4207 case Intrinsic::spv_udot:
4208 case Intrinsic::spv_sdot:
4209 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4211 return selectIntegerDot(ResVReg, ResType,
I,
4212 IID == Intrinsic::spv_sdot);
4213 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4214 case Intrinsic::spv_dot4add_i8packed:
4215 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4217 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4218 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4219 case Intrinsic::spv_dot4add_u8packed:
4220 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4222 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4223 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4224 case Intrinsic::spv_all:
4225 return selectAll(ResVReg, ResType,
I);
4226 case Intrinsic::spv_any:
4227 return selectAny(ResVReg, ResType,
I);
4228 case Intrinsic::spv_cross:
4229 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4230 case Intrinsic::spv_distance:
4231 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4232 case Intrinsic::spv_lerp:
4233 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4234 case Intrinsic::spv_length:
4235 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4236 case Intrinsic::spv_degrees:
4237 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4238 case Intrinsic::spv_faceforward:
4239 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4240 case Intrinsic::spv_frac:
4241 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4242 case Intrinsic::spv_isinf:
4243 return selectOpIsInf(ResVReg, ResType,
I);
4244 case Intrinsic::spv_isnan:
4245 return selectOpIsNan(ResVReg, ResType,
I);
4246 case Intrinsic::spv_normalize:
4247 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4248 case Intrinsic::spv_refract:
4249 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4250 case Intrinsic::spv_reflect:
4251 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4252 case Intrinsic::spv_rsqrt:
4253 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4254 case Intrinsic::spv_sign:
4255 return selectSign(ResVReg, ResType,
I);
4256 case Intrinsic::spv_smoothstep:
4257 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4258 case Intrinsic::spv_firstbituhigh:
4259 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4260 case Intrinsic::spv_firstbitshigh:
4261 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4262 case Intrinsic::spv_firstbitlow:
4263 return selectFirstBitLow(ResVReg, ResType,
I);
4264 case Intrinsic::spv_group_memory_barrier_with_group_sync: {
4266 buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent,
I);
4267 Register ScopeReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
4269 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpControlBarrier))
4276 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4277 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4278 SPIRV::StorageClass::StorageClass ResSC =
4282 "Generic storage class");
4283 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4291 case Intrinsic::spv_lifetime_start:
4292 case Intrinsic::spv_lifetime_end: {
4293 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
4294 : SPIRV::OpLifetimeStop;
4295 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
4296 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
4305 case Intrinsic::spv_saturate:
4306 return selectSaturate(ResVReg, ResType,
I);
4307 case Intrinsic::spv_nclamp:
4308 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
4309 case Intrinsic::spv_uclamp:
4310 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
4311 case Intrinsic::spv_sclamp:
4312 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
4313 case Intrinsic::spv_subgroup_prefix_bit_count:
4314 return selectWavePrefixBitCount(ResVReg, ResType,
I);
4315 case Intrinsic::spv_wave_active_countbits:
4316 return selectWaveActiveCountBits(ResVReg, ResType,
I);
4317 case Intrinsic::spv_wave_all_equal:
4318 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
4319 case Intrinsic::spv_wave_all:
4320 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
4321 case Intrinsic::spv_wave_any:
4322 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
4323 case Intrinsic::spv_subgroup_ballot:
4324 return selectWaveOpInst(ResVReg, ResType,
I,
4325 SPIRV::OpGroupNonUniformBallot);
4326 case Intrinsic::spv_wave_is_first_lane:
4327 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
4328 case Intrinsic::spv_wave_reduce_or:
4329 return selectWaveReduceOp(ResVReg, ResType,
I,
4330 SPIRV::OpGroupNonUniformBitwiseOr);
4331 case Intrinsic::spv_wave_reduce_umax:
4332 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
4333 case Intrinsic::spv_wave_reduce_max:
4334 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
4335 case Intrinsic::spv_wave_reduce_umin:
4336 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
4337 case Intrinsic::spv_wave_reduce_min:
4338 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
4339 case Intrinsic::spv_wave_reduce_sum:
4340 return selectWaveReduceSum(ResVReg, ResType,
I);
4341 case Intrinsic::spv_wave_product:
4342 return selectWaveReduceProduct(ResVReg, ResType,
I);
4343 case Intrinsic::spv_wave_readlane:
4344 return selectWaveOpInst(ResVReg, ResType,
I,
4345 SPIRV::OpGroupNonUniformShuffle);
4346 case Intrinsic::spv_wave_prefix_sum:
4347 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
4348 case Intrinsic::spv_wave_prefix_product:
4349 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
4350 case Intrinsic::spv_step:
4351 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
4352 case Intrinsic::spv_radians:
4353 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
4357 case Intrinsic::instrprof_increment:
4358 case Intrinsic::instrprof_increment_step:
4359 case Intrinsic::instrprof_value_profile:
4362 case Intrinsic::spv_value_md:
4364 case Intrinsic::spv_resource_handlefrombinding: {
4365 return selectHandleFromBinding(ResVReg, ResType,
I);
4367 case Intrinsic::spv_resource_counterhandlefrombinding:
4368 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
4369 case Intrinsic::spv_resource_updatecounter:
4370 return selectUpdateCounter(ResVReg, ResType,
I);
4371 case Intrinsic::spv_resource_store_typedbuffer: {
4372 return selectImageWriteIntrinsic(
I);
4374 case Intrinsic::spv_resource_load_typedbuffer: {
4375 return selectReadImageIntrinsic(ResVReg, ResType,
I);
4377 case Intrinsic::spv_resource_sample:
4378 case Intrinsic::spv_resource_sample_clamp:
4379 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
4380 case Intrinsic::spv_resource_samplebias:
4381 case Intrinsic::spv_resource_samplebias_clamp:
4382 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
4383 case Intrinsic::spv_resource_samplegrad:
4384 case Intrinsic::spv_resource_samplegrad_clamp:
4385 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
4386 case Intrinsic::spv_resource_samplelevel:
4387 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
4388 case Intrinsic::spv_resource_samplecmp:
4389 case Intrinsic::spv_resource_samplecmp_clamp:
4390 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
4391 case Intrinsic::spv_resource_samplecmplevelzero:
4392 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
4393 case Intrinsic::spv_resource_gather:
4394 case Intrinsic::spv_resource_gather_cmp:
4395 return selectGatherIntrinsic(ResVReg, ResType,
I);
4396 case Intrinsic::spv_resource_getpointer: {
4397 return selectResourceGetPointer(ResVReg, ResType,
I);
4399 case Intrinsic::spv_pushconstant_getpointer: {
4400 return selectPushConstantGetPointer(ResVReg, ResType,
I);
4402 case Intrinsic::spv_discard: {
4403 return selectDiscard(ResVReg, ResType,
I);
4405 case Intrinsic::spv_resource_nonuniformindex: {
4406 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
4408 case Intrinsic::spv_unpackhalf2x16: {
4409 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
4411 case Intrinsic::spv_packhalf2x16: {
4412 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
4414 case Intrinsic::spv_ddx:
4415 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
4416 case Intrinsic::spv_ddy:
4417 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
4418 case Intrinsic::spv_ddx_coarse:
4419 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
4420 case Intrinsic::spv_ddy_coarse:
4421 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
4422 case Intrinsic::spv_ddx_fine:
4423 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
4424 case Intrinsic::spv_ddy_fine:
4425 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
4426 case Intrinsic::spv_fwidth:
4427 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
4428 case Intrinsic::spv_masked_gather:
4429 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4430 return selectMaskedGather(ResVReg, ResType,
I);
4431 return diagnoseUnsupported(
4432 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
4433 case Intrinsic::spv_masked_scatter:
4434 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4435 return selectMaskedScatter(
I);
4436 return diagnoseUnsupported(
4437 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
4439 std::string DiagMsg;
4440 raw_string_ostream OS(DiagMsg);
4442 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
4449bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
4450 SPIRVTypeInst ResType,
4451 MachineInstr &
I)
const {
4454 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
4461bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
4462 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4464 assert(Intr.getIntrinsicID() ==
4465 Intrinsic::spv_resource_counterhandlefrombinding);
4468 Register MainHandleReg = Intr.getOperand(2).getReg();
4470 assert(MainHandleDef->getIntrinsicID() ==
4471 Intrinsic::spv_resource_handlefrombinding);
4475 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
4476 Register IndexReg = MainHandleDef->getOperand(5).getReg();
4477 std::string CounterName =
4482 MachineIRBuilder MIRBuilder(
I);
4484 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
4486 ArraySize, IndexReg, CounterName, MIRBuilder);
4488 return BuildCOPY(ResVReg, CounterVarReg,
I);
4491bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
4492 SPIRVTypeInst ResType,
4493 MachineInstr &
I)
const {
4495 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
4497 Register CounterHandleReg = Intr.getOperand(2).getReg();
4498 Register IncrReg = Intr.getOperand(3).getReg();
4505 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
4506 assert(CounterVarPointeeType &&
4507 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
4508 "Counter variable must be a struct");
4510 SPIRV::StorageClass::StorageBuffer &&
4511 "Counter variable must be in the storage buffer storage class");
4513 "Counter variable must have exactly 1 member in the struct");
4514 const SPIRVTypeInst MemberType =
4517 "Counter variable struct must have a single i32 member");
4521 MachineIRBuilder MIRBuilder(
I);
4523 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
4526 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
4532 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4535 .
addUse(CounterHandleReg)
4542 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
4545 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
4548 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
4557 return BuildCOPY(ResVReg, AtomicRes,
I);
4565 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
4573bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
4574 SPIRVTypeInst ResType,
4575 MachineInstr &
I)
const {
4583 Register ImageReg =
I.getOperand(2).getReg();
4591 Register IdxReg =
I.getOperand(3).getReg();
4593 MachineInstr &Pos =
I;
4595 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
4599bool SPIRVInstructionSelector::generateSampleImage(
4602 DebugLoc Loc, MachineInstr &Pos)
const {
4613 if (!loadHandleBeforePosition(NewSamplerReg,
4619 MachineIRBuilder MIRBuilder(Pos);
4632 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
4633 ImOps.Lod.has_value();
4634 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
4635 : SPIRV::OpImageSampleImplicitLod;
4637 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
4638 : SPIRV::OpImageSampleDrefImplicitLod;
4647 MIB.
addUse(*ImOps.Compare);
4649 uint32_t ImageOperands = 0;
4651 ImageOperands |= SPIRV::ImageOperand::Bias;
4653 ImageOperands |= SPIRV::ImageOperand::Lod;
4654 if (ImOps.GradX && ImOps.GradY)
4655 ImageOperands |= SPIRV::ImageOperand::Grad;
4656 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
4658 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4661 "Non-constant offsets are not supported in sample instructions.");
4665 ImageOperands |= SPIRV::ImageOperand::MinLod;
4667 if (ImageOperands != 0) {
4668 MIB.
addImm(ImageOperands);
4669 if (ImageOperands & SPIRV::ImageOperand::Bias)
4671 if (ImageOperands & SPIRV::ImageOperand::Lod)
4673 if (ImageOperands & SPIRV::ImageOperand::Grad) {
4674 MIB.
addUse(*ImOps.GradX);
4675 MIB.
addUse(*ImOps.GradY);
4678 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4679 MIB.
addUse(*ImOps.Offset);
4680 if (ImageOperands & SPIRV::ImageOperand::MinLod)
4681 MIB.
addUse(*ImOps.MinLod);
4688bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
4689 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4690 Register ImageReg =
I.getOperand(2).getReg();
4691 Register SamplerReg =
I.getOperand(3).getReg();
4692 Register CoordinateReg =
I.getOperand(4).getReg();
4693 ImageOperands ImOps;
4694 if (
I.getNumOperands() > 5)
4695 ImOps.Offset =
I.getOperand(5).getReg();
4696 if (
I.getNumOperands() > 6)
4697 ImOps.MinLod =
I.getOperand(6).getReg();
4698 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4699 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4702bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
4703 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4704 Register ImageReg =
I.getOperand(2).getReg();
4705 Register SamplerReg =
I.getOperand(3).getReg();
4706 Register CoordinateReg =
I.getOperand(4).getReg();
4707 ImageOperands ImOps;
4708 ImOps.Bias =
I.getOperand(5).getReg();
4709 if (
I.getNumOperands() > 6)
4710 ImOps.Offset =
I.getOperand(6).getReg();
4711 if (
I.getNumOperands() > 7)
4712 ImOps.MinLod =
I.getOperand(7).getReg();
4713 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4714 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4717bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
4718 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4719 Register ImageReg =
I.getOperand(2).getReg();
4720 Register SamplerReg =
I.getOperand(3).getReg();
4721 Register CoordinateReg =
I.getOperand(4).getReg();
4722 ImageOperands ImOps;
4723 ImOps.GradX =
I.getOperand(5).getReg();
4724 ImOps.GradY =
I.getOperand(6).getReg();
4725 if (
I.getNumOperands() > 7)
4726 ImOps.Offset =
I.getOperand(7).getReg();
4727 if (
I.getNumOperands() > 8)
4728 ImOps.MinLod =
I.getOperand(8).getReg();
4729 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4730 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4733bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
4734 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4735 Register ImageReg =
I.getOperand(2).getReg();
4736 Register SamplerReg =
I.getOperand(3).getReg();
4737 Register CoordinateReg =
I.getOperand(4).getReg();
4738 ImageOperands ImOps;
4739 ImOps.Lod =
I.getOperand(5).getReg();
4740 if (
I.getNumOperands() > 6)
4741 ImOps.Offset =
I.getOperand(6).getReg();
4742 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4743 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4746bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
4747 SPIRVTypeInst ResType,
4748 MachineInstr &
I)
const {
4749 Register ImageReg =
I.getOperand(2).getReg();
4750 Register SamplerReg =
I.getOperand(3).getReg();
4751 Register CoordinateReg =
I.getOperand(4).getReg();
4752 ImageOperands ImOps;
4753 ImOps.Compare =
I.getOperand(5).getReg();
4754 if (
I.getNumOperands() > 6)
4755 ImOps.Offset =
I.getOperand(6).getReg();
4756 if (
I.getNumOperands() > 7)
4757 ImOps.MinLod =
I.getOperand(7).getReg();
4758 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4759 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4762bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
4763 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4764 Register ImageReg =
I.getOperand(2).getReg();
4765 Register SamplerReg =
I.getOperand(3).getReg();
4766 Register CoordinateReg =
I.getOperand(4).getReg();
4767 ImageOperands ImOps;
4768 ImOps.Compare =
I.getOperand(5).getReg();
4769 if (
I.getNumOperands() > 6)
4770 ImOps.Offset =
I.getOperand(6).getReg();
4773 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4774 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4777bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
4778 SPIRVTypeInst ResType,
4779 MachineInstr &
I)
const {
4780 Register ImageReg =
I.getOperand(2).getReg();
4781 Register SamplerReg =
I.getOperand(3).getReg();
4782 Register CoordinateReg =
I.getOperand(4).getReg();
4785 "ImageReg is not an image type.");
4790 ComponentOrCompareReg =
I.getOperand(5).getReg();
4791 OffsetReg =
I.getOperand(6).getReg();
4794 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
4798 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
4799 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
4800 Dim != SPIRV::Dim::DIM_Rect) {
4802 "Gather operations are only supported for 2D, Cube, and Rect images.");
4809 if (!loadHandleBeforePosition(
4814 MachineIRBuilder MIRBuilder(
I);
4815 SPIRVTypeInst SampledImageType =
4820 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
4828 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
4830 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
4832 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4837 .
addUse(ComponentOrCompareReg);
4839 uint32_t ImageOperands = 0;
4840 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
4841 if (Dim == SPIRV::Dim::DIM_Cube) {
4843 "Gather operations with offset are not supported for Cube images.");
4847 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4849 ImageOperands |= SPIRV::ImageOperand::Offset;
4853 if (ImageOperands != 0) {
4854 MIB.
addImm(ImageOperands);
4856 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4864bool SPIRVInstructionSelector::generateImageReadOrFetch(
4869 "ImageReg is not an image type.");
4871 bool IsSignedInteger =
4876 bool IsFetch = (SampledOp.getImm() == 1);
4879 if (ResultSize == 4) {
4882 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
4888 if (IsSignedInteger)
4894 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
4898 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
4903 if (IsSignedInteger)
4907 if (ResultSize == 1) {
4916 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
4919bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
4920 SPIRVTypeInst ResType,
4921 MachineInstr &
I)
const {
4922 Register ResourcePtr =
I.getOperand(2).getReg();
4924 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
4933 MachineIRBuilder MIRBuilder(
I);
4935 Register IndexReg =
I.getOperand(3).getReg();
4938 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4948bool SPIRVInstructionSelector::selectPushConstantGetPointer(
4949 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4954bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
4955 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4956 Register ObjReg =
I.getOperand(2).getReg();
4957 if (!BuildCOPY(ResVReg, ObjReg,
I))
4967 decorateUsesAsNonUniform(ResVReg);
4971void SPIRVInstructionSelector::decorateUsesAsNonUniform(
4974 while (WorkList.
size() > 0) {
4978 bool IsDecorated =
false;
4980 if (
Use.getOpcode() == SPIRV::OpDecorate &&
4981 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
4987 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
4989 if (ResultReg == CurrentReg)
4997 SPIRV::Decoration::NonUniformEXT, {});
5002bool SPIRVInstructionSelector::extractSubvector(
5004 MachineInstr &InsertionPoint)
const {
5006 [[maybe_unused]] uint64_t InputSize =
5009 assert(InputSize > 1 &&
"The input must be a vector.");
5010 assert(ResultSize > 1 &&
"The result must be a vector.");
5011 assert(ResultSize < InputSize &&
5012 "Cannot extract more element than there are in the input.");
5015 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
5016 for (uint64_t
I = 0;
I < ResultSize;
I++) {
5019 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5028 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
5030 TII.get(SPIRV::OpCompositeConstruct))
5034 for (
Register ComponentReg : ComponentRegisters)
5035 MIB.
addUse(ComponentReg);
5040bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
5041 MachineInstr &
I)
const {
5048 Register ImageReg =
I.getOperand(1).getReg();
5056 Register CoordinateReg =
I.getOperand(2).getReg();
5057 Register DataReg =
I.getOperand(3).getReg();
5060 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
5068Register SPIRVInstructionSelector::buildPointerToResource(
5069 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
5070 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
5071 StringRef Name, MachineIRBuilder MIRBuilder)
const {
5073 if (ArraySize == 1) {
5074 SPIRVTypeInst PtrType =
5077 "SpirvResType did not have an explicit layout.");
5082 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
5083 SPIRVTypeInst VarPointerType =
5086 VarPointerType, Set,
Binding, Name, MIRBuilder);
5088 SPIRVTypeInst ResPointerType =
5101bool SPIRVInstructionSelector::selectFirstBitSet16(
5102 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5103 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
5105 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
5109 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
5112bool SPIRVInstructionSelector::selectFirstBitSet32(
5114 unsigned BitSetOpcode)
const {
5115 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5118 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5125bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
5127 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5134 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
5136 MachineIRBuilder MIRBuilder(
I);
5139 SPIRVTypeInst I64x2Type =
5141 SPIRVTypeInst Vec2ResType =
5144 std::vector<Register> PartialRegs;
5147 unsigned CurrentComponent = 0;
5148 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
5154 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5155 TII.get(SPIRV::OpVectorShuffle))
5160 .
addImm(CurrentComponent)
5161 .
addImm(CurrentComponent + 1);
5168 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType,
I, BitSetResult,
5169 BitSetOpcode, SwapPrimarySide))
5172 PartialRegs.push_back(SubVecBitSetReg);
5176 if (CurrentComponent != ComponentCount) {
5182 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
5183 SPIRV::OpVectorExtractDynamic))
5189 if (!selectFirstBitSet64(FinalElemBitSetReg,
BaseType,
I, FinalElemReg,
5190 BitSetOpcode, SwapPrimarySide))
5193 PartialRegs.push_back(FinalElemBitSetReg);
5198 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
5199 SPIRV::OpCompositeConstruct);
5202bool SPIRVInstructionSelector::selectFirstBitSet64(
5204 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5217 if (ComponentCount > 2) {
5218 return selectFirstBitSet64Overflow(ResVReg, ResType,
I, SrcReg,
5219 BitSetOpcode, SwapPrimarySide);
5223 MachineIRBuilder MIRBuilder(
I);
5225 BaseType, 2 * ComponentCount, MIRBuilder,
false);
5229 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
5235 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
5242 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
5245 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntZero},
5246 SPIRV::OpVectorExtractDynamic))
5248 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntOne},
5249 SPIRV::OpVectorExtractDynamic))
5253 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5254 TII.get(SPIRV::OpVectorShuffle))
5262 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
5268 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5269 TII.get(SPIRV::OpVectorShuffle))
5277 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
5297 SelectOp = SPIRV::OpSelectSISCond;
5298 AddOp = SPIRV::OpIAddS;
5306 SelectOp = SPIRV::OpSelectVIVCond;
5307 AddOp = SPIRV::OpIAddV;
5317 if (SwapPrimarySide) {
5318 PrimaryReg = LowReg;
5319 SecondaryReg = HighReg;
5320 PrimaryShiftReg = Reg0;
5321 SecondaryShiftReg = Reg32;
5326 if (!selectOpWithSrcs(BReg, BoolType,
I, {PrimaryReg, NegOneReg},
5332 if (!selectOpWithSrcs(TmpReg, ResType,
I, {
BReg, SecondaryReg, PrimaryReg},
5338 if (!selectOpWithSrcs(ValReg, ResType,
I,
5339 {
BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))
5342 return selectOpWithSrcs(ResVReg, ResType,
I, {ValReg, TmpReg}, AddOp);
5345bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
5346 SPIRVTypeInst ResType,
5348 bool IsSigned)
const {
5350 Register OpReg =
I.getOperand(2).getReg();
5353 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5354 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
5358 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5360 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5362 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5366 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
5370bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
5371 SPIRVTypeInst ResType,
5372 MachineInstr &
I)
const {
5374 Register OpReg =
I.getOperand(2).getReg();
5379 unsigned ExtendOpcode = SPIRV::OpUConvert;
5380 unsigned BitSetOpcode = GL::FindILsb;
5384 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5386 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5388 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5395bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
5396 SPIRVTypeInst ResType,
5397 MachineInstr &
I)
const {
5401 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
5404 .
addUse(
I.getOperand(2).getReg())
5407 unsigned Alignment =
I.getOperand(3).getImm();
5413bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
5414 SPIRVTypeInst ResType,
5415 MachineInstr &
I)
const {
5419 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
5422 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
5425 unsigned Alignment =
I.getOperand(2).getImm();
5432bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
5437 const MachineInstr *PrevI =
I.getPrevNode();
5439 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
5443 .
addMBB(
I.getOperand(0).getMBB())
5448 .
addMBB(
I.getOperand(0).getMBB())
5453bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
5464 const MachineInstr *NextI =
I.getNextNode();
5466 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
5472 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
5474 .
addUse(
I.getOperand(0).getReg())
5475 .
addMBB(
I.getOperand(1).getMBB())
5481bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
5482 MachineInstr &
I)
const {
5484 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
5486 const unsigned NumOps =
I.getNumOperands();
5487 for (
unsigned i = 1; i <
NumOps; i += 2) {
5488 MIB.
addUse(
I.getOperand(i + 0).getReg());
5489 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
5495bool SPIRVInstructionSelector::selectGlobalValue(
5496 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
5498 MachineIRBuilder MIRBuilder(
I);
5499 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
5502 std::string GlobalIdent;
5504 unsigned &
ID = UnnamedGlobalIDs[GV];
5506 ID = UnnamedGlobalIDs.
size();
5507 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
5533 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
5540 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
5545 MachineInstrBuilder MIB1 =
5546 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5549 MachineInstrBuilder MIB2 =
5551 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
5555 GR.
add(ConstVal, MIB2);
5563 MachineInstrBuilder MIB3 =
5564 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5567 GR.
add(ConstVal, MIB3);
5571 assert(NewReg != ResVReg);
5572 return BuildCOPY(ResVReg, NewReg,
I);
5582 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
5588 SPIRVTypeInst ResType =
5592 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
5597 if (
GlobalVar->isExternallyInitialized() &&
5598 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
5599 constexpr unsigned ReadWriteINTEL = 3u;
5602 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
5608bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
5609 SPIRVTypeInst ResType,
5610 MachineInstr &
I)
const {
5612 return selectExtInst(ResVReg, ResType,
I, CL::log10);
5620 MachineIRBuilder MIRBuilder(
I);
5625 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5628 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5630 .
add(
I.getOperand(1))
5635 ResType->
getOpcode() == SPIRV::OpTypeFloat);
5637 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
5645 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
5646 ? SPIRV::OpVectorTimesScalar
5657bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
5658 SPIRVTypeInst ResType,
5659 MachineInstr &
I)
const {
5662 return selectExtInst(ResVReg, ResType,
I, CL::pown);
5668 Register ExpReg =
I.getOperand(2).getReg();
5670 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
5671 SPIRV::OpConvertSToF))
5673 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
5680bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
5681 SPIRVTypeInst ResType,
5682 MachineInstr &
I)
const {
5698 MachineIRBuilder MIRBuilder(
I);
5701 ResType, MIRBuilder, SPIRV::StorageClass::Function);
5713 MachineBasicBlock &EntryBB =
I.getMF()->front();
5717 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
5720 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
5726 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5729 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
5732 .
add(
I.getOperand(
I.getNumExplicitDefs()))
5736 Register IntegralPartReg =
I.getOperand(1).getReg();
5737 if (IntegralPartReg.
isValid()) {
5739 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5750 assert(
false &&
"GLSL::Modf is deprecated.");
5761bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
5762 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5763 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5764 MachineIRBuilder MIRBuilder(
I);
5765 const SPIRVTypeInst Vec3Ty =
5768 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
5780 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5784 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
5790 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5797 assert(
I.getOperand(2).isReg());
5798 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
5802 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5813bool SPIRVInstructionSelector::loadBuiltinInputID(
5814 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5815 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5816 MachineIRBuilder MIRBuilder(
I);
5818 ResType, MIRBuilder, SPIRV::StorageClass::Input);
5833 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5837 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5846SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
5847 MachineInstr &
I)
const {
5848 MachineIRBuilder MIRBuilder(
I);
5849 if (
Type->getOpcode() != SPIRV::OpTypeVector)
5853 if (VectorSize == 4)
5861bool SPIRVInstructionSelector::loadHandleBeforePosition(
5862 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
5863 MachineInstr &Pos)
const {
5866 Intrinsic::spv_resource_handlefrombinding);
5874 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
5875 MachineIRBuilder MIRBuilder(HandleDef);
5876 SPIRVTypeInst VarType = ResType;
5877 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
5879 if (IsStructuredBuffer) {
5885 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
5886 ArraySize, IndexReg, Name, MIRBuilder);
5890 uint32_t LoadOpcode =
5891 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
5901void SPIRVInstructionSelector::errorIfInstrOutsideShader(
5902 MachineInstr &
I)
const {
5904 std::string DiagMsg;
5905 raw_string_ostream OS(DiagMsg);
5906 I.print(OS,
true,
false,
false,
false);
5907 DiagMsg +=
" is only supported in shaders.\n";
5913InstructionSelector *
5917 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
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 TableGen::Emitter::OptClass< SkeletonEmitter > X("gen-skeleton-class", "Generate example skeleton class")
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.
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...