32#include "llvm/IR/IntrinsicsSPIRV.h"
36#define DEBUG_TYPE "spirv-isel"
43 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
48 std::optional<Register> Bias;
49 std::optional<Register>
Offset;
50 std::optional<Register> MinLod;
51 std::optional<Register> GradX;
52 std::optional<Register> GradY;
53 std::optional<Register> Lod;
54 std::optional<Register> Compare;
57llvm::SPIRV::SelectionControl::SelectionControl
58getSelectionOperandForImm(
int Imm) {
60 return SPIRV::SelectionControl::Flatten;
62 return SPIRV::SelectionControl::DontFlatten;
64 return SPIRV::SelectionControl::None;
68#define GET_GLOBALISEL_PREDICATE_BITSET
69#include "SPIRVGenGlobalISel.inc"
70#undef GET_GLOBALISEL_PREDICATE_BITSET
97#define GET_GLOBALISEL_PREDICATES_DECL
98#include "SPIRVGenGlobalISel.inc"
99#undef GET_GLOBALISEL_PREDICATES_DECL
101#define GET_GLOBALISEL_TEMPORARIES_DECL
102#include "SPIRVGenGlobalISel.inc"
103#undef GET_GLOBALISEL_TEMPORARIES_DECL
127 unsigned BitSetOpcode)
const;
131 unsigned BitSetOpcode)
const;
135 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
139 unsigned BitSetOpcode,
140 bool SwapPrimarySide)
const;
147 unsigned Opcode)
const;
150 unsigned Opcode)
const;
169 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
180 unsigned OpType)
const;
229 template <
bool Signed>
232 template <
bool Signed>
239 template <
typename PickOpcodeFn>
242 PickOpcodeFn &&PickOpcode)
const;
253 template <
typename PickOpcodeFn>
256 PickOpcodeFn &&PickOpcode)
const;
271 bool IsSigned)
const;
273 bool IsSigned,
unsigned Opcode)
const;
275 bool IsSigned)
const;
281 bool IsSigned)
const;
314 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
315 bool useMISrc =
true,
317 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
318 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
319 bool useMISrc =
true,
321 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
322 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
323 bool setMIFlags =
true,
bool useMISrc =
true,
325 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
326 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
327 bool useMISrc =
true,
330 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
331 MachineInstr &
I)
const;
333 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
334 MachineInstr &
I)
const;
336 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
337 MachineInstr &
I,
unsigned Opcode)
const;
339 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
340 MachineInstr &
I)
const;
342 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
343 MachineInstr &
I)
const;
347 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
348 MachineInstr &
I)
const;
350 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
351 MachineInstr &
I)
const;
353 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
354 MachineInstr &
I)
const;
355 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
356 MachineInstr &
I)
const;
357 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
358 MachineInstr &
I)
const;
359 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
360 MachineInstr &
I)
const;
361 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
362 MachineInstr &
I)
const;
363 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
364 MachineInstr &
I)
const;
365 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
366 SPIRVTypeInst ResType,
367 MachineInstr &
I)
const;
368 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
369 MachineInstr &
I)
const;
370 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
371 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
372 MachineInstr &
I)
const;
373 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
374 MachineInstr &
I)
const;
375 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
376 MachineInstr &
I)
const;
377 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
378 MachineInstr &
I)
const;
379 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
380 MachineInstr &
I)
const;
381 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
382 MachineInstr &
I)
const;
383 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
384 MachineInstr &
I)
const;
385 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
386 MachineInstr &
I)
const;
387 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
388 MachineInstr &
I,
const unsigned DPdOpCode)
const;
390 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
391 SPIRVTypeInst ResType =
nullptr)
const;
393 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
394 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
395 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
397 MachineInstr &
I)
const;
398 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
400 bool wrapIntoSpecConstantOp(MachineInstr &
I,
403 Register getUcharPtrTypeReg(MachineInstr &
I,
404 SPIRV::StorageClass::StorageClass SC)
const;
405 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
407 uint32_t Opcode)
const;
408 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
409 SPIRVTypeInst SrcPtrTy)
const;
410 Register buildPointerToResource(SPIRVTypeInst ResType,
411 SPIRV::StorageClass::StorageClass SC,
412 uint32_t Set, uint32_t
Binding,
413 uint32_t ArraySize,
Register IndexReg,
415 MachineIRBuilder MIRBuilder)
const;
416 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
417 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
418 Register &ReadReg, MachineInstr &InsertionPoint)
const;
419 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
421 DebugLoc Loc, MachineInstr &Pos)
const;
422 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
424 Register CoordinateReg,
const ImageOperands &ImOps,
427 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
428 Register ResVReg, SPIRVTypeInst ResType,
429 MachineInstr &
I)
const;
430 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
431 Register ResVReg, SPIRVTypeInst ResType,
432 MachineInstr &
I)
const;
433 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
434 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
435 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
436 void errorIfInstrOutsideShader(MachineInstr &
I)
const;
439bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
441 if (
TET->getTargetExtName() ==
"spirv.Image") {
444 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
445 return TET->getTypeParameter(0)->isIntegerTy();
449#define GET_GLOBALISEL_IMPL
450#include "SPIRVGenGlobalISel.inc"
451#undef GET_GLOBALISEL_IMPL
457 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
460#include
"SPIRVGenGlobalISel.inc"
463#include
"SPIRVGenGlobalISel.inc"
475 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
479void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
480 if (HasVRegsReset == &MF)
485 for (
unsigned I = 0,
E =
MRI.getNumVirtRegs();
I !=
E; ++
I) {
495 for (
const auto &
MBB : MF) {
496 for (
const auto &
MI :
MBB) {
499 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
503 LLT DstType =
MRI.getType(DstReg);
505 LLT SrcType =
MRI.getType(SrcReg);
506 if (DstType != SrcType)
507 MRI.setType(DstReg,
MRI.getType(SrcReg));
509 const TargetRegisterClass *DstRC =
MRI.getRegClassOrNull(DstReg);
510 const TargetRegisterClass *SrcRC =
MRI.getRegClassOrNull(SrcReg);
511 if (DstRC != SrcRC && SrcRC)
512 MRI.setRegClass(DstReg, SrcRC);
523 while (!Stack.empty()) {
528 switch (
MI->getOpcode()) {
529 case TargetOpcode::G_INTRINSIC:
530 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
531 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
533 Intrinsic::spv_const_composite)
536 case TargetOpcode::G_BUILD_VECTOR:
537 case TargetOpcode::G_SPLAT_VECTOR:
539 i < OpDef->getNumOperands(); i++) {
544 Stack.push_back(OpNestedDef);
547 case TargetOpcode::G_CONSTANT:
548 case TargetOpcode::G_FCONSTANT:
549 case TargetOpcode::G_IMPLICIT_DEF:
550 case SPIRV::OpConstantTrue:
551 case SPIRV::OpConstantFalse:
552 case SPIRV::OpConstantI:
553 case SPIRV::OpConstantF:
554 case SPIRV::OpConstantComposite:
555 case SPIRV::OpConstantCompositeContinuedINTEL:
556 case SPIRV::OpConstantSampler:
557 case SPIRV::OpConstantNull:
559 case SPIRV::OpConstantFunctionPointerINTEL:
586 case Intrinsic::spv_all:
587 case Intrinsic::spv_alloca:
588 case Intrinsic::spv_any:
589 case Intrinsic::spv_bitcast:
590 case Intrinsic::spv_const_composite:
591 case Intrinsic::spv_cross:
592 case Intrinsic::spv_degrees:
593 case Intrinsic::spv_distance:
594 case Intrinsic::spv_extractelt:
595 case Intrinsic::spv_extractv:
596 case Intrinsic::spv_faceforward:
597 case Intrinsic::spv_fdot:
598 case Intrinsic::spv_firstbitlow:
599 case Intrinsic::spv_firstbitshigh:
600 case Intrinsic::spv_firstbituhigh:
601 case Intrinsic::spv_frac:
602 case Intrinsic::spv_gep:
603 case Intrinsic::spv_global_offset:
604 case Intrinsic::spv_global_size:
605 case Intrinsic::spv_group_id:
606 case Intrinsic::spv_insertelt:
607 case Intrinsic::spv_insertv:
608 case Intrinsic::spv_isinf:
609 case Intrinsic::spv_isnan:
610 case Intrinsic::spv_lerp:
611 case Intrinsic::spv_length:
612 case Intrinsic::spv_normalize:
613 case Intrinsic::spv_num_subgroups:
614 case Intrinsic::spv_num_workgroups:
615 case Intrinsic::spv_ptrcast:
616 case Intrinsic::spv_radians:
617 case Intrinsic::spv_reflect:
618 case Intrinsic::spv_refract:
619 case Intrinsic::spv_resource_getpointer:
620 case Intrinsic::spv_resource_handlefrombinding:
621 case Intrinsic::spv_resource_handlefromimplicitbinding:
622 case Intrinsic::spv_resource_nonuniformindex:
623 case Intrinsic::spv_resource_sample:
624 case Intrinsic::spv_rsqrt:
625 case Intrinsic::spv_saturate:
626 case Intrinsic::spv_sdot:
627 case Intrinsic::spv_sign:
628 case Intrinsic::spv_smoothstep:
629 case Intrinsic::spv_step:
630 case Intrinsic::spv_subgroup_id:
631 case Intrinsic::spv_subgroup_local_invocation_id:
632 case Intrinsic::spv_subgroup_max_size:
633 case Intrinsic::spv_subgroup_size:
634 case Intrinsic::spv_thread_id:
635 case Intrinsic::spv_thread_id_in_group:
636 case Intrinsic::spv_udot:
637 case Intrinsic::spv_undef:
638 case Intrinsic::spv_value_md:
639 case Intrinsic::spv_workgroup_size:
651 case SPIRV::OpTypeVoid:
652 case SPIRV::OpTypeBool:
653 case SPIRV::OpTypeInt:
654 case SPIRV::OpTypeFloat:
655 case SPIRV::OpTypeVector:
656 case SPIRV::OpTypeMatrix:
657 case SPIRV::OpTypeImage:
658 case SPIRV::OpTypeSampler:
659 case SPIRV::OpTypeSampledImage:
660 case SPIRV::OpTypeArray:
661 case SPIRV::OpTypeRuntimeArray:
662 case SPIRV::OpTypeStruct:
663 case SPIRV::OpTypeOpaque:
664 case SPIRV::OpTypePointer:
665 case SPIRV::OpTypeFunction:
666 case SPIRV::OpTypeEvent:
667 case SPIRV::OpTypeDeviceEvent:
668 case SPIRV::OpTypeReserveId:
669 case SPIRV::OpTypeQueue:
670 case SPIRV::OpTypePipe:
671 case SPIRV::OpTypeForwardPointer:
672 case SPIRV::OpTypePipeStorage:
673 case SPIRV::OpTypeNamedBarrier:
674 case SPIRV::OpTypeAccelerationStructureNV:
675 case SPIRV::OpTypeCooperativeMatrixNV:
676 case SPIRV::OpTypeCooperativeMatrixKHR:
686 if (
MI.getNumDefs() == 0)
689 for (
const auto &MO :
MI.all_defs()) {
691 if (
Reg.isPhysical()) {
695 for (
const auto &
UseMI :
MRI.use_nodbg_instructions(
Reg)) {
696 if (
UseMI.getOpcode() != SPIRV::OpName) {
703 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
704 MI.isLifetimeMarker()) {
707 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
718 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
719 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
722 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
727 if (
MI.mayStore() ||
MI.isCall() ||
728 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
729 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
730 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
741 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
748void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
750 for (
const auto &MO :
MI.all_defs()) {
754 SmallVector<MachineInstr *, 4> UselessOpNames;
755 for (MachineInstr &
UseMI :
MRI->use_nodbg_instructions(
Reg)) {
757 "There is still a use of the dead function.");
760 for (MachineInstr *OpNameMI : UselessOpNames) {
762 OpNameMI->eraseFromParent();
767void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
770 removeOpNamesForDeadMI(
MI);
771 MI.eraseFromParent();
774bool SPIRVInstructionSelector::select(MachineInstr &
I) {
775 resetVRegsType(*
I.getParent()->getParent());
777 assert(
I.getParent() &&
"Instruction should be in a basic block!");
778 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
783 removeDeadInstruction(
I);
790 if (Opcode == SPIRV::ASSIGN_TYPE) {
791 Register DstReg =
I.getOperand(0).getReg();
792 Register SrcReg =
I.getOperand(1).getReg();
793 auto *
Def =
MRI->getVRegDef(SrcReg);
795 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
796 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
797 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
798 Register SelectDstReg =
Def->getOperand(0).getReg();
799 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
801 assert(SuccessToSelectSelect);
803 Def->eraseFromParent();
804 MRI->replaceRegWith(DstReg, SelectDstReg);
810 bool Res = selectImpl(
I, *CoverageInfo);
812 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
813 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
817 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
824 MRI->setRegClass(SrcReg,
MRI->getRegClass(DstReg));
825 MRI->replaceRegWith(SrcReg, DstReg);
829 }
else if (
I.getNumDefs() == 1) {
841 removeDeadInstruction(
I);
846 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
847 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
853 bool HasDefs =
I.getNumDefs() > 0;
856 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
857 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
858 if (spvSelect(ResVReg, ResType,
I)) {
860 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
871 case TargetOpcode::G_CONSTANT:
872 case TargetOpcode::G_FCONSTANT:
874 case TargetOpcode::G_SADDO:
875 case TargetOpcode::G_SSUBO:
882 MachineInstr &
I)
const {
883 const TargetRegisterClass *DstRC =
MRI->getRegClassOrNull(DestReg);
884 const TargetRegisterClass *SrcRC =
MRI->getRegClassOrNull(SrcReg);
885 if (DstRC != SrcRC && SrcRC)
886 MRI->setRegClass(DestReg, SrcRC);
887 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
894bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
895 SPIRVTypeInst ResType,
896 MachineInstr &
I)
const {
897 const unsigned Opcode =
I.getOpcode();
899 return selectImpl(
I, *CoverageInfo);
901 case TargetOpcode::G_CONSTANT:
902 case TargetOpcode::G_FCONSTANT:
903 return selectConst(ResVReg, ResType,
I);
904 case TargetOpcode::G_GLOBAL_VALUE:
905 return selectGlobalValue(ResVReg,
I);
906 case TargetOpcode::G_IMPLICIT_DEF:
907 return selectOpUndef(ResVReg, ResType,
I);
908 case TargetOpcode::G_FREEZE:
909 return selectFreeze(ResVReg, ResType,
I);
911 case TargetOpcode::G_INTRINSIC:
912 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
913 case TargetOpcode::G_INTRINSIC_CONVERGENT:
914 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
915 return selectIntrinsic(ResVReg, ResType,
I);
916 case TargetOpcode::G_BITREVERSE:
917 return selectBitreverse(ResVReg, ResType,
I);
919 case TargetOpcode::G_BUILD_VECTOR:
920 return selectBuildVector(ResVReg, ResType,
I);
921 case TargetOpcode::G_SPLAT_VECTOR:
922 return selectSplatVector(ResVReg, ResType,
I);
924 case TargetOpcode::G_SHUFFLE_VECTOR: {
925 MachineBasicBlock &BB = *
I.getParent();
926 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
929 .
addUse(
I.getOperand(1).getReg())
930 .
addUse(
I.getOperand(2).getReg());
931 for (
auto V :
I.getOperand(3).getShuffleMask())
936 case TargetOpcode::G_MEMMOVE:
937 case TargetOpcode::G_MEMCPY:
938 case TargetOpcode::G_MEMSET:
939 return selectMemOperation(ResVReg,
I);
941 case TargetOpcode::G_ICMP:
942 return selectICmp(ResVReg, ResType,
I);
943 case TargetOpcode::G_FCMP:
944 return selectFCmp(ResVReg, ResType,
I);
946 case TargetOpcode::G_FRAME_INDEX:
947 return selectFrameIndex(ResVReg, ResType,
I);
949 case TargetOpcode::G_LOAD:
950 return selectLoad(ResVReg, ResType,
I);
951 case TargetOpcode::G_STORE:
952 return selectStore(
I);
954 case TargetOpcode::G_BR:
955 return selectBranch(
I);
956 case TargetOpcode::G_BRCOND:
957 return selectBranchCond(
I);
959 case TargetOpcode::G_PHI:
960 return selectPhi(ResVReg,
I);
962 case TargetOpcode::G_FPTOSI:
963 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
964 case TargetOpcode::G_FPTOUI:
965 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
967 case TargetOpcode::G_FPTOSI_SAT:
968 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
969 case TargetOpcode::G_FPTOUI_SAT:
970 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
972 case TargetOpcode::G_SITOFP:
973 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
974 case TargetOpcode::G_UITOFP:
975 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
977 case TargetOpcode::G_CTPOP:
978 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitCount);
979 case TargetOpcode::G_SMIN:
980 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
981 case TargetOpcode::G_UMIN:
982 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
984 case TargetOpcode::G_SMAX:
985 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
986 case TargetOpcode::G_UMAX:
987 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
989 case TargetOpcode::G_SCMP:
990 return selectSUCmp(ResVReg, ResType,
I,
true);
991 case TargetOpcode::G_UCMP:
992 return selectSUCmp(ResVReg, ResType,
I,
false);
993 case TargetOpcode::G_LROUND:
994 case TargetOpcode::G_LLROUND: {
996 MRI->createVirtualRegister(
MRI->getRegClass(ResVReg),
"lround");
997 MRI->setRegClass(regForLround, &SPIRV::iIDRegClass);
999 regForLround, *(
I.getParent()->getParent()));
1001 CL::round, GL::Round,
false);
1003 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1010 case TargetOpcode::G_STRICT_FMA:
1011 case TargetOpcode::G_FMA: {
1014 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1017 .
addUse(
I.getOperand(1).getReg())
1018 .
addUse(
I.getOperand(2).getReg())
1019 .
addUse(
I.getOperand(3).getReg())
1024 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1027 case TargetOpcode::G_STRICT_FLDEXP:
1028 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1030 case TargetOpcode::G_FPOW:
1031 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1032 case TargetOpcode::G_FPOWI:
1033 return selectExtInst(ResVReg, ResType,
I, CL::pown);
1035 case TargetOpcode::G_FEXP:
1036 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1037 case TargetOpcode::G_FEXP2:
1038 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1039 case TargetOpcode::G_FEXP10:
1040 return selectExp10(ResVReg, ResType,
I);
1042 case TargetOpcode::G_FMODF:
1043 return selectModf(ResVReg, ResType,
I);
1044 case TargetOpcode::G_FSINCOS:
1045 return selectSincos(ResVReg, ResType,
I);
1047 case TargetOpcode::G_FLOG:
1048 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1049 case TargetOpcode::G_FLOG2:
1050 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1051 case TargetOpcode::G_FLOG10:
1052 return selectLog10(ResVReg, ResType,
I);
1054 case TargetOpcode::G_FABS:
1055 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1056 case TargetOpcode::G_ABS:
1057 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1059 case TargetOpcode::G_FMINNUM:
1060 case TargetOpcode::G_FMINIMUM:
1061 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1062 case TargetOpcode::G_FMAXNUM:
1063 case TargetOpcode::G_FMAXIMUM:
1064 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1066 case TargetOpcode::G_FCOPYSIGN:
1067 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1069 case TargetOpcode::G_FCEIL:
1070 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1071 case TargetOpcode::G_FFLOOR:
1072 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1074 case TargetOpcode::G_FCOS:
1075 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1076 case TargetOpcode::G_FSIN:
1077 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1078 case TargetOpcode::G_FTAN:
1079 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1080 case TargetOpcode::G_FACOS:
1081 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1082 case TargetOpcode::G_FASIN:
1083 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1084 case TargetOpcode::G_FATAN:
1085 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1086 case TargetOpcode::G_FATAN2:
1087 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1088 case TargetOpcode::G_FCOSH:
1089 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1090 case TargetOpcode::G_FSINH:
1091 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1092 case TargetOpcode::G_FTANH:
1093 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1095 case TargetOpcode::G_STRICT_FSQRT:
1096 case TargetOpcode::G_FSQRT:
1097 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1099 case TargetOpcode::G_CTTZ:
1100 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1101 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1102 case TargetOpcode::G_CTLZ:
1103 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1104 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1106 case TargetOpcode::G_INTRINSIC_ROUND:
1107 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1108 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1109 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1110 case TargetOpcode::G_INTRINSIC_TRUNC:
1111 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1112 case TargetOpcode::G_FRINT:
1113 case TargetOpcode::G_FNEARBYINT:
1114 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1116 case TargetOpcode::G_SMULH:
1117 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1118 case TargetOpcode::G_UMULH:
1119 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1121 case TargetOpcode::G_SADDSAT:
1122 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1123 case TargetOpcode::G_UADDSAT:
1124 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1125 case TargetOpcode::G_SSUBSAT:
1126 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1127 case TargetOpcode::G_USUBSAT:
1128 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1130 case TargetOpcode::G_FFREXP:
1131 return selectFrexp(ResVReg, ResType,
I);
1133 case TargetOpcode::G_UADDO:
1134 return selectOverflowArith(ResVReg, ResType,
I,
1135 ResType->
getOpcode() == SPIRV::OpTypeVector
1136 ? SPIRV::OpIAddCarryV
1137 : SPIRV::OpIAddCarryS);
1138 case TargetOpcode::G_USUBO:
1139 return selectOverflowArith(ResVReg, ResType,
I,
1140 ResType->
getOpcode() == SPIRV::OpTypeVector
1141 ? SPIRV::OpISubBorrowV
1142 : SPIRV::OpISubBorrowS);
1143 case TargetOpcode::G_UMULO:
1144 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1145 case TargetOpcode::G_SMULO:
1146 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1148 case TargetOpcode::G_SEXT:
1149 return selectExt(ResVReg, ResType,
I,
true);
1150 case TargetOpcode::G_ANYEXT:
1151 case TargetOpcode::G_ZEXT:
1152 return selectExt(ResVReg, ResType,
I,
false);
1153 case TargetOpcode::G_TRUNC:
1154 return selectTrunc(ResVReg, ResType,
I);
1155 case TargetOpcode::G_FPTRUNC:
1156 case TargetOpcode::G_FPEXT:
1157 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1159 case TargetOpcode::G_PTRTOINT:
1160 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1161 case TargetOpcode::G_INTTOPTR:
1162 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1163 case TargetOpcode::G_BITCAST:
1164 return selectBitcast(ResVReg, ResType,
I);
1165 case TargetOpcode::G_ADDRSPACE_CAST:
1166 return selectAddrSpaceCast(ResVReg, ResType,
I);
1167 case TargetOpcode::G_PTR_ADD: {
1169 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1173 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1174 (*II).getOpcode() == TargetOpcode::COPY ||
1175 (*II).getOpcode() == SPIRV::OpVariable) &&
1178 bool IsGVInit =
false;
1180 UseIt =
MRI->use_instr_begin(
I.getOperand(0).getReg()),
1181 UseEnd =
MRI->use_instr_end();
1182 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1183 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1184 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1185 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1195 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1198 Register NewVReg =
MRI->createGenericVirtualRegister(
MRI->getType(GV));
1199 MRI->setRegClass(NewVReg,
MRI->getRegClass(GV));
1208 "incompatible result and operand types in a bitcast");
1210 MachineInstrBuilder MIB =
1211 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1218 : SPIRV::OpInBoundsPtrAccessChain))
1222 .
addUse(
I.getOperand(2).getReg())
1225 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1229 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1231 .
addUse(
I.getOperand(2).getReg())
1240 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1243 .
addImm(
static_cast<uint32_t
>(
1244 SPIRV::Opcode::InBoundsPtrAccessChain))
1247 .
addUse(
I.getOperand(2).getReg());
1252 case TargetOpcode::G_ATOMICRMW_OR:
1253 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1254 case TargetOpcode::G_ATOMICRMW_ADD:
1255 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1256 case TargetOpcode::G_ATOMICRMW_AND:
1257 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1258 case TargetOpcode::G_ATOMICRMW_MAX:
1259 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1260 case TargetOpcode::G_ATOMICRMW_MIN:
1261 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1262 case TargetOpcode::G_ATOMICRMW_SUB:
1263 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1264 case TargetOpcode::G_ATOMICRMW_XOR:
1265 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1266 case TargetOpcode::G_ATOMICRMW_UMAX:
1267 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1268 case TargetOpcode::G_ATOMICRMW_UMIN:
1269 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1270 case TargetOpcode::G_ATOMICRMW_XCHG:
1271 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1272 case TargetOpcode::G_ATOMIC_CMPXCHG:
1273 return selectAtomicCmpXchg(ResVReg, ResType,
I);
1275 case TargetOpcode::G_ATOMICRMW_FADD:
1276 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1277 case TargetOpcode::G_ATOMICRMW_FSUB:
1279 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1280 ResType->
getOpcode() == SPIRV::OpTypeVector
1282 : SPIRV::OpFNegate);
1283 case TargetOpcode::G_ATOMICRMW_FMIN:
1284 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1285 case TargetOpcode::G_ATOMICRMW_FMAX:
1286 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1288 case TargetOpcode::G_FENCE:
1289 return selectFence(
I);
1291 case TargetOpcode::G_STACKSAVE:
1292 return selectStackSave(ResVReg, ResType,
I);
1293 case TargetOpcode::G_STACKRESTORE:
1294 return selectStackRestore(
I);
1296 case TargetOpcode::G_UNMERGE_VALUES:
1302 case TargetOpcode::G_TRAP:
1303 case TargetOpcode::G_UBSANTRAP:
1304 case TargetOpcode::DBG_LABEL:
1306 case TargetOpcode::G_DEBUGTRAP:
1307 return selectDebugTrap(ResVReg, ResType,
I);
1314bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1315 SPIRVTypeInst ResType,
1316 MachineInstr &
I)
const {
1317 unsigned Opcode = SPIRV::OpNop;
1324bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1325 SPIRVTypeInst ResType,
1327 GL::GLSLExtInst GLInst,
1328 bool setMIFlags,
bool useMISrc,
1331 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1332 std::string DiagMsg;
1333 raw_string_ostream OS(DiagMsg);
1334 I.print(OS,
true,
false,
false,
false);
1335 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1338 return selectExtInst(ResVReg, ResType,
I,
1339 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1340 setMIFlags, useMISrc, SrcRegs);
1343bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1344 SPIRVTypeInst ResType,
1346 CL::OpenCLExtInst CLInst,
1347 bool setMIFlags,
bool useMISrc,
1349 return selectExtInst(ResVReg, ResType,
I,
1350 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1351 setMIFlags, useMISrc, SrcRegs);
1354bool SPIRVInstructionSelector::selectExtInst(
1355 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1356 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1358 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1359 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1360 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1364bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1365 SPIRVTypeInst ResType,
1368 bool setMIFlags,
bool useMISrc,
1371 for (
const auto &[InstructionSet, Opcode] : Insts) {
1375 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1378 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1383 const unsigned NumOps =
I.getNumOperands();
1386 I.getOperand(Index).getType() ==
1387 MachineOperand::MachineOperandType::MO_IntrinsicID)
1390 MIB.
add(
I.getOperand(Index));
1402bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1403 SPIRVTypeInst ResType,
1404 MachineInstr &
I)
const {
1405 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1406 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1407 for (
const auto &Ex : ExtInsts) {
1408 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1409 uint32_t Opcode = Ex.second;
1413 MachineIRBuilder MIRBuilder(
I);
1416 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1421 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1424 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1427 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1430 .
addImm(
static_cast<uint32_t
>(Ex.first))
1432 .
add(
I.getOperand(2))
1436 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1437 .
addDef(
I.getOperand(1).getReg())
1446bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1447 SPIRVTypeInst ResType,
1448 MachineInstr &
I)
const {
1449 Register CosResVReg =
I.getOperand(1).getReg();
1450 unsigned SrcIdx =
I.getNumExplicitDefs();
1455 MachineIRBuilder MIRBuilder(
I);
1457 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1462 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1465 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1467 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1470 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1472 .
add(
I.getOperand(SrcIdx))
1475 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1483 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1486 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1488 .
add(
I.getOperand(SrcIdx))
1490 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1493 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1495 .
add(
I.getOperand(SrcIdx))
1502bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1503 SPIRVTypeInst ResType,
1505 std::vector<Register> Srcs,
1506 unsigned Opcode)
const {
1507 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1517bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1518 SPIRVTypeInst ResType,
1520 unsigned Opcode)
const {
1522 Register SrcReg =
I.getOperand(1).getReg();
1525 MRI->def_instr_begin(SrcReg);
1526 DefIt !=
MRI->def_instr_end(); DefIt = std::next(DefIt)) {
1527 unsigned DefOpCode = DefIt->getOpcode();
1528 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1531 if (
auto *VRD =
getVRegDef(*
MRI, DefIt->getOperand(1).getReg()))
1532 DefOpCode = VRD->getOpcode();
1534 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1535 DefOpCode == TargetOpcode::G_CONSTANT ||
1536 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1542 uint32_t SpecOpcode = 0;
1544 case SPIRV::OpConvertPtrToU:
1545 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1547 case SPIRV::OpConvertUToPtr:
1548 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1553 TII.get(SPIRV::OpSpecConstantOp))
1563 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1567bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1568 SPIRVTypeInst ResType,
1569 MachineInstr &
I)
const {
1570 Register OpReg =
I.getOperand(1).getReg();
1571 SPIRVTypeInst OpType =
1575 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1585 if (
MemOp->isVolatile())
1586 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1587 if (
MemOp->isNonTemporal())
1588 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1590 if (!ST->isShader() &&
MemOp->getAlign().value())
1591 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1595 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1596 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1600 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1602 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1606 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1610 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1612 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1624 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1626 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1628 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1632bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1633 SPIRVTypeInst ResType,
1634 MachineInstr &
I)
const {
1636 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1641 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1642 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1644 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1646 MRI->createVirtualRegister(
MRI->getRegClass(HandleReg));
1648 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1652 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1653 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1654 I.getDebugLoc(),
I);
1658 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1662 if (!
I.getNumMemOperands()) {
1663 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1665 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1668 MachineIRBuilder MIRBuilder(
I);
1675bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1677 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1678 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1683 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1684 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1686 MRI->createVirtualRegister(
MRI->getRegClass(HandleReg));
1689 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1693 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1694 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1695 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1696 TII.get(SPIRV::OpImageWrite))
1702 if (sampledTypeIsSignedInteger(LLVMHandleType))
1705 BMI.constrainAllUses(
TII,
TRI, RBI);
1711 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpStore))
1714 if (!
I.getNumMemOperands()) {
1715 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1717 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1720 MachineIRBuilder MIRBuilder(
I);
1727bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
1728 SPIRVTypeInst ResType,
1729 MachineInstr &
I)
const {
1730 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1732 "llvm.stacksave intrinsic: this instruction requires the following "
1733 "SPIR-V extension: SPV_INTEL_variable_length_array",
1736 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
1743bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
1744 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1746 "llvm.stackrestore intrinsic: this instruction requires the following "
1747 "SPIR-V extension: SPV_INTEL_variable_length_array",
1749 if (!
I.getOperand(0).isReg())
1752 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
1753 .
addUse(
I.getOperand(0).getReg())
1759SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
1760 MachineIRBuilder MIRBuilder(
I);
1761 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1768 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
1772 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1773 Type *ArrTy = ArrayType::get(ValTy, Num);
1775 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
1778 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
1785 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
1788 .
addImm(SPIRV::StorageClass::UniformConstant)
1799bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
1802 Register DstReg =
I.getOperand(0).getReg();
1812 "Unable to determine pointee type size for OpCopyMemory");
1813 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
1814 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
1816 "OpCopyMemory requires the size to match the pointee type size");
1817 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
1820 if (
I.getNumMemOperands()) {
1821 MachineIRBuilder MIRBuilder(
I);
1828bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
1831 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
1832 .
addUse(
I.getOperand(0).getReg())
1834 .
addUse(
I.getOperand(2).getReg());
1835 if (
I.getNumMemOperands()) {
1836 MachineIRBuilder MIRBuilder(
I);
1843bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
1844 MachineInstr &
I)
const {
1845 Register SrcReg =
I.getOperand(1).getReg();
1846 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
1847 Register VarReg = getOrCreateMemSetGlobal(
I);
1850 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1852 ValTy,
I, SPIRV::StorageClass::UniformConstant);
1854 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
1858 if (!selectCopyMemory(
I, SrcReg))
1861 if (!selectCopyMemorySized(
I, SrcReg))
1864 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
1865 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
1870bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
1871 SPIRVTypeInst ResType,
1874 unsigned NegateOpcode)
const {
1876 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1879 Register ScopeReg = buildI32Constant(Scope,
I);
1881 Register Ptr =
I.getOperand(1).getReg();
1887 Register MemSemReg = buildI32Constant(MemSem ,
I);
1889 Register ValueReg =
I.getOperand(2).getReg();
1890 if (NegateOpcode != 0) {
1893 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
1898 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
1909bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
1910 unsigned ArgI =
I.getNumOperands() - 1;
1912 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
1913 SPIRVTypeInst SrcType =
1915 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
1917 "cannot select G_UNMERGE_VALUES with a non-vector argument");
1919 SPIRVTypeInst ScalarType =
1922 unsigned CurrentIndex = 0;
1923 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
1924 Register ResVReg =
I.getOperand(i).getReg();
1927 LLT ResLLT =
MRI->getType(ResVReg);
1933 ResType = ScalarType;
1939 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
1942 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1948 for (
unsigned j = 0;
j < NumElements; ++
j) {
1949 MIB.
addImm(CurrentIndex + j);
1951 CurrentIndex += NumElements;
1955 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
1967bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
1970 Register MemSemReg = buildI32Constant(MemSem,
I);
1972 uint32_t
Scope =
static_cast<uint32_t
>(
1974 Register ScopeReg = buildI32Constant(Scope,
I);
1976 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
1983bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
1984 SPIRVTypeInst ResType,
1986 unsigned Opcode)
const {
1987 Type *ResTy =
nullptr;
1991 "Not enough info to select the arithmetic with overflow instruction");
1994 "with overflow instruction");
2000 MachineIRBuilder MIRBuilder(
I);
2002 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2003 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2009 Register ZeroReg = buildZerosVal(ResType,
I);
2012 MRI->setRegClass(StructVReg, &SPIRV::IDRegClass);
2014 if (ResName.
size() > 0)
2019 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2022 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2023 MIB.
addUse(
I.getOperand(i).getReg());
2028 MRI->setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2029 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2031 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2032 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2039 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2040 .
addDef(
I.getOperand(1).getReg())
2048bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2049 SPIRVTypeInst ResType,
2050 MachineInstr &
I)
const {
2054 Register Ptr =
I.getOperand(2).getReg();
2057 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2060 ScopeReg = buildI32Constant(Scope,
I);
2062 unsigned ScSem =
static_cast<uint32_t
>(
2065 unsigned MemSemEq =
static_cast<uint32_t
>(
getMemSemantics(AO)) | ScSem;
2066 Register MemSemEqReg = buildI32Constant(MemSemEq,
I);
2068 unsigned MemSemNeq =
static_cast<uint32_t
>(
getMemSemantics(FO)) | ScSem;
2069 if (MemSemEq == MemSemNeq)
2070 MemSemNeqReg = MemSemEqReg;
2072 MemSemNeqReg = buildI32Constant(MemSemEq,
I);
2075 ScopeReg =
I.getOperand(5).getReg();
2076 MemSemEqReg =
I.getOperand(6).getReg();
2077 MemSemNeqReg =
I.getOperand(7).getReg();
2081 Register Val =
I.getOperand(4).getReg();
2085 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2104 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2111 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2123 case SPIRV::StorageClass::DeviceOnlyINTEL:
2124 case SPIRV::StorageClass::HostOnlyINTEL:
2133 bool IsGRef =
false;
2134 bool IsAllowedRefs =
2135 llvm::all_of(
MRI->use_instructions(ResVReg), [&IsGRef](
auto const &It) {
2136 unsigned Opcode = It.getOpcode();
2137 if (Opcode == SPIRV::OpConstantComposite ||
2138 Opcode == SPIRV::OpVariable ||
2139 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2140 return IsGRef = true;
2141 return Opcode == SPIRV::OpName;
2143 return IsAllowedRefs && IsGRef;
2146Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2147 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2149 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2153SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2155 uint32_t Opcode)
const {
2156 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2157 TII.get(SPIRV::OpSpecConstantOp))
2165SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2166 SPIRVTypeInst SrcPtrTy)
const {
2167 SPIRVTypeInst GenericPtrTy =
2169 Register Tmp =
MRI->createVirtualRegister(&SPIRV::pIDRegClass);
2171 SPIRV::StorageClass::Generic),
2173 MachineFunction *MF =
I.getParent()->getParent();
2175 MachineInstrBuilder MIB = buildSpecConstantOp(
2177 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2187bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2188 SPIRVTypeInst ResType,
2189 MachineInstr &
I)
const {
2193 Register SrcPtr =
I.getOperand(1).getReg();
2197 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2198 ResType->
getOpcode() != SPIRV::OpTypePointer)
2199 return BuildCOPY(ResVReg, SrcPtr,
I);
2209 unsigned SpecOpcode =
2211 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2214 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2221 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2223 .constrainAllUses(
TII,
TRI, RBI);
2225 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2227 buildSpecConstantOp(
2229 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2230 .constrainAllUses(
TII,
TRI, RBI);
2237 return BuildCOPY(ResVReg, SrcPtr,
I);
2239 if ((SrcSC == SPIRV::StorageClass::Function &&
2240 DstSC == SPIRV::StorageClass::Private) ||
2241 (DstSC == SPIRV::StorageClass::Function &&
2242 SrcSC == SPIRV::StorageClass::Private))
2243 return BuildCOPY(ResVReg, SrcPtr,
I);
2247 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2250 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2253 SPIRVTypeInst GenericPtrTy =
2272 return selectUnOp(ResVReg, ResType,
I,
2273 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2275 return selectUnOp(ResVReg, ResType,
I,
2276 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2278 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2280 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2290 return SPIRV::OpFOrdEqual;
2292 return SPIRV::OpFOrdGreaterThanEqual;
2294 return SPIRV::OpFOrdGreaterThan;
2296 return SPIRV::OpFOrdLessThanEqual;
2298 return SPIRV::OpFOrdLessThan;
2300 return SPIRV::OpFOrdNotEqual;
2302 return SPIRV::OpOrdered;
2304 return SPIRV::OpFUnordEqual;
2306 return SPIRV::OpFUnordGreaterThanEqual;
2308 return SPIRV::OpFUnordGreaterThan;
2310 return SPIRV::OpFUnordLessThanEqual;
2312 return SPIRV::OpFUnordLessThan;
2314 return SPIRV::OpFUnordNotEqual;
2316 return SPIRV::OpUnordered;
2326 return SPIRV::OpIEqual;
2328 return SPIRV::OpINotEqual;
2330 return SPIRV::OpSGreaterThanEqual;
2332 return SPIRV::OpSGreaterThan;
2334 return SPIRV::OpSLessThanEqual;
2336 return SPIRV::OpSLessThan;
2338 return SPIRV::OpUGreaterThanEqual;
2340 return SPIRV::OpUGreaterThan;
2342 return SPIRV::OpULessThanEqual;
2344 return SPIRV::OpULessThan;
2353 return SPIRV::OpPtrEqual;
2355 return SPIRV::OpPtrNotEqual;
2366 return SPIRV::OpLogicalEqual;
2368 return SPIRV::OpLogicalNotEqual;
2402bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2403 SPIRVTypeInst ResType,
2405 unsigned OpAnyOrAll)
const {
2406 assert(
I.getNumOperands() == 3);
2407 assert(
I.getOperand(2).isReg());
2409 Register InputRegister =
I.getOperand(2).getReg();
2416 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2417 if (IsBoolTy && !IsVectorTy) {
2418 assert(ResVReg ==
I.getOperand(0).getReg());
2419 return BuildCOPY(ResVReg, InputRegister,
I);
2423 unsigned SpirvNotEqualId =
2424 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2426 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2431 IsBoolTy ? InputRegister
2439 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2441 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2458bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2459 SPIRVTypeInst ResType,
2460 MachineInstr &
I)
const {
2461 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2464bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2465 SPIRVTypeInst ResType,
2466 MachineInstr &
I)
const {
2467 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2471bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2472 SPIRVTypeInst ResType,
2473 MachineInstr &
I)
const {
2474 assert(
I.getNumOperands() == 4);
2475 assert(
I.getOperand(2).isReg());
2476 assert(
I.getOperand(3).isReg());
2478 [[maybe_unused]] SPIRVTypeInst VecType =
2483 "dot product requires a vector of at least 2 components");
2485 [[maybe_unused]] SPIRVTypeInst EltType =
2494 .
addUse(
I.getOperand(2).getReg())
2495 .
addUse(
I.getOperand(3).getReg())
2500bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2501 SPIRVTypeInst ResType,
2504 assert(
I.getNumOperands() == 4);
2505 assert(
I.getOperand(2).isReg());
2506 assert(
I.getOperand(3).isReg());
2509 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2513 .
addUse(
I.getOperand(2).getReg())
2514 .
addUse(
I.getOperand(3).getReg())
2521bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2522 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2523 assert(
I.getNumOperands() == 4);
2524 assert(
I.getOperand(2).isReg());
2525 assert(
I.getOperand(3).isReg());
2529 Register Vec0 =
I.getOperand(2).getReg();
2530 Register Vec1 =
I.getOperand(3).getReg();
2534 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2543 "dot product requires a vector of at least 2 components");
2546 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2556 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2567 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2579bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2580 SPIRVTypeInst ResType,
2581 MachineInstr &
I)
const {
2583 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2586 .
addUse(
I.getOperand(2).getReg())
2591bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2592 SPIRVTypeInst ResType,
2593 MachineInstr &
I)
const {
2595 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2598 .
addUse(
I.getOperand(2).getReg())
2603template <
bool Signed>
2604bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2605 SPIRVTypeInst ResType,
2606 MachineInstr &
I)
const {
2607 assert(
I.getNumOperands() == 5);
2608 assert(
I.getOperand(2).isReg());
2609 assert(
I.getOperand(3).isReg());
2610 assert(
I.getOperand(4).isReg());
2613 Register Acc =
I.getOperand(2).getReg();
2617 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2619 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
2624 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
2627 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2639template <
bool Signed>
2640bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
2641 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2642 assert(
I.getNumOperands() == 5);
2643 assert(
I.getOperand(2).isReg());
2644 assert(
I.getOperand(3).isReg());
2645 assert(
I.getOperand(4).isReg());
2648 Register Acc =
I.getOperand(2).getReg();
2654 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
2658 for (
unsigned i = 0; i < 4; i++) {
2660 Register AElt =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2670 Register BElt =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2681 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
2689 Register MaskMul =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2700 i < 3 ?
MRI->createVirtualRegister(&SPIRV::IDRegClass) : ResVReg;
2701 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2716bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
2717 SPIRVTypeInst ResType,
2718 MachineInstr &
I)
const {
2719 assert(
I.getNumOperands() == 3);
2720 assert(
I.getOperand(2).isReg());
2722 Register VZero = buildZerosValF(ResType,
I);
2723 Register VOne = buildOnesValF(ResType,
I);
2725 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
2728 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2730 .
addUse(
I.getOperand(2).getReg())
2737bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
2738 SPIRVTypeInst ResType,
2739 MachineInstr &
I)
const {
2740 assert(
I.getNumOperands() == 3);
2741 assert(
I.getOperand(2).isReg());
2743 Register InputRegister =
I.getOperand(2).getReg();
2745 auto &
DL =
I.getDebugLoc();
2755 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
2757 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
2759 ?
MRI->createVirtualRegister(&SPIRV::IDRegClass)
2765 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2770 if (NeedsConversion) {
2771 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
2782bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
2783 SPIRVTypeInst ResType,
2785 unsigned Opcode)
const {
2789 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2795 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
2796 BMI.addUse(
I.getOperand(J).getReg());
2803bool SPIRVInstructionSelector::selectWaveActiveCountBits(
2804 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2809 if (!selectWaveOpInst(BallotReg, BallotType,
I,
2810 SPIRV::OpGroupNonUniformBallot))
2815 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2820 .
addImm(SPIRV::GroupOperation::Reduce)
2829 if (
Type->getOpcode() != SPIRV::OpTypeVector)
2833 return Type->getOperand(2).getImm();
2836bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
2837 SPIRVTypeInst ResType,
2838 MachineInstr &
I)
const {
2843 Register InputReg =
I.getOperand(2).getReg();
2848 bool IsVector = NumElems > 1;
2851 SPIRVTypeInst ElemInputType = InputType;
2852 SPIRVTypeInst ElemBoolType = ResType;
2865 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
2866 SPIRV::OpGroupNonUniformAllEqual);
2871 ElementResults.
reserve(NumElems);
2873 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
2886 ElemInput = Extracted;
2892 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
2903 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
2914bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
2915 SPIRVTypeInst ResType,
2916 MachineInstr &
I)
const {
2918 assert(
I.getNumOperands() == 3);
2920 auto Op =
I.getOperand(2);
2932 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
2943 Register BallotVReg =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2954 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2958 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
2965bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
2966 SPIRVTypeInst ResType,
2968 bool IsUnsigned)
const {
2969 return selectWaveReduce(
2970 ResVReg, ResType,
I, IsUnsigned,
2971 [&](
Register InputRegister,
bool IsUnsigned) {
2972 const bool IsFloatTy =
2974 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
2975 : SPIRV::OpGroupNonUniformSMax;
2976 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
2980bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
2981 SPIRVTypeInst ResType,
2983 bool IsUnsigned)
const {
2984 return selectWaveReduce(
2985 ResVReg, ResType,
I, IsUnsigned,
2986 [&](
Register InputRegister,
bool IsUnsigned) {
2987 const bool IsFloatTy =
2989 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
2990 : SPIRV::OpGroupNonUniformSMin;
2991 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
2995bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
2996 SPIRVTypeInst ResType,
2997 MachineInstr &
I)
const {
2998 return selectWaveReduce(ResVReg, ResType,
I,
false,
2999 [&](
Register InputRegister,
bool IsUnsigned) {
3001 InputRegister, SPIRV::OpTypeFloat);
3002 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3003 : SPIRV::OpGroupNonUniformIAdd;
3007template <
typename PickOpcodeFn>
3008bool SPIRVInstructionSelector::selectWaveReduce(
3009 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3010 PickOpcodeFn &&PickOpcode)
const {
3011 assert(
I.getNumOperands() == 3);
3012 assert(
I.getOperand(2).isReg());
3014 Register InputRegister =
I.getOperand(2).getReg();
3021 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3027 .
addImm(SPIRV::GroupOperation::Reduce)
3028 .
addUse(
I.getOperand(2).getReg())
3033bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3034 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3035 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3036 [&](
Register InputRegister,
bool IsUnsigned) {
3038 InputRegister, SPIRV::OpTypeFloat);
3040 ? SPIRV::OpGroupNonUniformFAdd
3041 : SPIRV::OpGroupNonUniformIAdd;
3045bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3046 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3047 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3048 [&](
Register InputRegister,
bool IsUnsigned) {
3050 InputRegister, SPIRV::OpTypeFloat);
3052 ? SPIRV::OpGroupNonUniformFMul
3053 : SPIRV::OpGroupNonUniformIMul;
3057template <
typename PickOpcodeFn>
3058bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3059 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3060 PickOpcodeFn &&PickOpcode)
const {
3061 assert(
I.getNumOperands() == 3);
3062 assert(
I.getOperand(2).isReg());
3064 Register InputRegister =
I.getOperand(2).getReg();
3071 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3077 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3078 .
addUse(
I.getOperand(2).getReg())
3083bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3084 SPIRVTypeInst ResType,
3085 MachineInstr &
I)
const {
3087 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3090 .
addUse(
I.getOperand(1).getReg())
3095bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3096 SPIRVTypeInst ResType,
3097 MachineInstr &
I)
const {
3103 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3105 Register OpReg =
I.getOperand(1).getReg();
3106 if (MachineInstr *Def =
MRI->getVRegDef(OpReg)) {
3107 if (
Def->getOpcode() == TargetOpcode::COPY)
3108 Def =
MRI->getVRegDef(
Def->getOperand(1).getReg());
3110 switch (
Def->getOpcode()) {
3111 case SPIRV::ASSIGN_TYPE:
3112 if (MachineInstr *AssignToDef =
3113 MRI->getVRegDef(
Def->getOperand(1).getReg())) {
3114 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3115 Reg =
Def->getOperand(2).getReg();
3118 case SPIRV::OpUndef:
3119 Reg =
Def->getOperand(1).getReg();
3122 unsigned DestOpCode;
3124 DestOpCode = SPIRV::OpConstantNull;
3126 DestOpCode = TargetOpcode::COPY;
3129 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3130 .
addDef(
I.getOperand(0).getReg())
3138bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3139 SPIRVTypeInst ResType,
3140 MachineInstr &
I)
const {
3142 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3144 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3148 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3153 for (
unsigned i =
I.getNumExplicitDefs();
3154 i <
I.getNumExplicitOperands() && IsConst; ++i)
3158 if (!IsConst &&
N < 2)
3160 "There must be at least two constituent operands in a vector");
3163 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3164 TII.get(IsConst ? SPIRV::OpConstantComposite
3165 : SPIRV::OpCompositeConstruct))
3168 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3169 MIB.
addUse(
I.getOperand(i).getReg());
3174bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3175 SPIRVTypeInst ResType,
3176 MachineInstr &
I)
const {
3178 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3180 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3186 if (!
I.getOperand(
OpIdx).isReg())
3193 if (!IsConst &&
N < 2)
3195 "There must be at least two constituent operands in a vector");
3198 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3199 TII.get(IsConst ? SPIRV::OpConstantComposite
3200 : SPIRV::OpCompositeConstruct))
3203 for (
unsigned i = 0; i <
N; ++i)
3209bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3210 SPIRVTypeInst ResType,
3211 MachineInstr &
I)
const {
3216 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3218 Opcode = SPIRV::OpDemoteToHelperInvocation;
3220 Opcode = SPIRV::OpKill;
3222 if (MachineInstr *NextI =
I.getNextNode()) {
3224 NextI->eraseFromParent();
3234bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3235 SPIRVTypeInst ResType,
unsigned CmpOpc,
3236 MachineInstr &
I)
const {
3237 Register Cmp0 =
I.getOperand(2).getReg();
3238 Register Cmp1 =
I.getOperand(3).getReg();
3241 "CMP operands should have the same type");
3242 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3252bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3253 SPIRVTypeInst ResType,
3254 MachineInstr &
I)
const {
3255 auto Pred =
I.getOperand(1).getPredicate();
3258 Register CmpOperand =
I.getOperand(2).getReg();
3265 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3269SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3270 SPIRVTypeInst ResType)
const {
3272 SPIRVTypeInst SpvI32Ty =
3275 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3282 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3285 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3288 .
addImm(APInt(32, Val).getZExtValue());
3290 GR.
add(ConstInt,
MI);
3295bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
3296 SPIRVTypeInst ResType,
3297 MachineInstr &
I)
const {
3299 return selectCmp(ResVReg, ResType, CmpOp,
I);
3302bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
3303 SPIRVTypeInst ResType,
3304 MachineInstr &
I)
const {
3306 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
3313 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
3314 ResType->
getOpcode() != SPIRV::OpTypeFloat)
3317 MachineIRBuilder MIRBuilder(
I);
3319 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
3325 "only float operands supported by GLSL extended math");
3328 MIRBuilder, SpirvScalarType);
3330 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
3331 ? SPIRV::OpVectorTimesScalar
3334 if (!selectOpWithSrcs(ArgReg, ResType,
I,
3335 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
3337 if (!selectExtInst(ResVReg, ResType,
I,
3338 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
3348Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
3349 MachineInstr &
I)
const {
3352 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3357bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
3363 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
3371 if (
Def->getOpcode() == SPIRV::OpConstantNull)
3374 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
3375 Def->getOpcode() == SPIRV::OpConstantI)
3384 MachineInstr *
Def =
MRI->getVRegDef(
Reg);
3388 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
3389 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
3391 Intrinsic::spv_const_composite)) {
3392 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
3393 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
3394 if (!IsZero(
Def->getOperand(i).getReg()))
3403Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
3404 MachineInstr &
I)
const {
3408 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3413Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
3414 MachineInstr &
I)
const {
3418 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3424 SPIRVTypeInst ResType,
3425 MachineInstr &
I)
const {
3429 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3434bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
3435 SPIRVTypeInst ResType,
3436 MachineInstr &
I)
const {
3437 Register SelectFirstArg =
I.getOperand(2).getReg();
3438 Register SelectSecondArg =
I.getOperand(3).getReg();
3447 SPIRV::OpTypeVector;
3454 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
3455 }
else if (IsPtrTy) {
3456 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
3458 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
3462 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
3463 }
else if (IsPtrTy) {
3464 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
3466 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3469 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3472 .
addUse(
I.getOperand(1).getReg())
3481bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
3482 SPIRVTypeInst ResType,
3484 MachineInstr &InsertAt,
3485 bool IsSigned)
const {
3487 Register ZeroReg = buildZerosVal(ResType, InsertAt);
3488 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
3489 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
3491 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3503bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
3504 SPIRVTypeInst ResType,
3505 MachineInstr &
I,
bool IsSigned,
3506 unsigned Opcode)
const {
3507 Register SrcReg =
I.getOperand(1).getReg();
3513 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
3518 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
3520 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
3523bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
3524 SPIRVTypeInst ResType, MachineInstr &
I,
3525 bool IsSigned)
const {
3526 Register SrcReg =
I.getOperand(1).getReg();
3528 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
3532 if (ResType == SrcType)
3533 return BuildCOPY(ResVReg, SrcReg,
I);
3535 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3536 return selectUnOp(ResVReg, ResType,
I, Opcode);
3539bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
3540 SPIRVTypeInst ResType,
3542 bool IsSigned)
const {
3543 MachineIRBuilder MIRBuilder(
I);
3544 MachineRegisterInfo *
MRI = MIRBuilder.getMRI();
3559 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
3562 .
addUse(
I.getOperand(1).getReg())
3563 .
addUse(
I.getOperand(2).getReg())
3569 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
3572 .
addUse(
I.getOperand(1).getReg())
3573 .
addUse(
I.getOperand(2).getReg())
3581 unsigned SelectOpcode =
3582 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
3587 .
addUse(buildOnesVal(
true, ResType,
I))
3588 .
addUse(buildZerosVal(ResType,
I))
3595 .
addUse(buildOnesVal(
false, ResType,
I))
3600bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
3603 SPIRVTypeInst IntTy,
3604 SPIRVTypeInst BoolTy)
const {
3607 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
3608 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
3610 Register One = buildOnesVal(
false, IntTy,
I);
3618 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
3627bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
3628 SPIRVTypeInst ResType,
3629 MachineInstr &
I)
const {
3630 Register IntReg =
I.getOperand(1).getReg();
3633 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
3634 if (ArgType == ResType)
3635 return BuildCOPY(ResVReg, IntReg,
I);
3637 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3638 return selectUnOp(ResVReg, ResType,
I, Opcode);
3641bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
3642 SPIRVTypeInst ResType,
3643 MachineInstr &
I)
const {
3644 unsigned Opcode =
I.getOpcode();
3645 unsigned TpOpcode = ResType->
getOpcode();
3647 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
3648 assert(Opcode == TargetOpcode::G_CONSTANT &&
3649 I.getOperand(1).getCImm()->isZero());
3650 MachineBasicBlock &DepMBB =
I.getMF()->front();
3653 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
3660 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
3663bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
3664 SPIRVTypeInst ResType,
3665 MachineInstr &
I)
const {
3666 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3673bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
3674 SPIRVTypeInst ResType,
3675 MachineInstr &
I)
const {
3677 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
3681 .
addUse(
I.getOperand(3).getReg())
3683 .
addUse(
I.getOperand(2).getReg());
3684 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
3690bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
3691 SPIRVTypeInst ResType,
3692 MachineInstr &
I)
const {
3693 Type *MaybeResTy =
nullptr;
3698 "Expected aggregate type for extractv instruction");
3700 SPIRV::AccessQualifier::ReadWrite,
false);
3704 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3707 .
addUse(
I.getOperand(2).getReg());
3708 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
3714bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
3715 SPIRVTypeInst ResType,
3716 MachineInstr &
I)
const {
3718 return selectInsertVal(ResVReg, ResType,
I);
3720 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
3723 .
addUse(
I.getOperand(2).getReg())
3724 .
addUse(
I.getOperand(3).getReg())
3725 .
addUse(
I.getOperand(4).getReg())
3730bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
3731 SPIRVTypeInst ResType,
3732 MachineInstr &
I)
const {
3734 return selectExtractVal(ResVReg, ResType,
I);
3736 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
3739 .
addUse(
I.getOperand(2).getReg())
3740 .
addUse(
I.getOperand(3).getReg())
3745bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
3746 SPIRVTypeInst ResType,
3747 MachineInstr &
I)
const {
3748 const bool IsGEPInBounds =
I.getOperand(2).getImm();
3754 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
3755 : SPIRV::OpAccessChain)
3756 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
3757 :
SPIRV::OpPtrAccessChain);
3759 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3763 .
addUse(
I.getOperand(3).getReg());
3765 (Opcode == SPIRV::OpPtrAccessChain ||
3766 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
3768 "Cannot translate GEP to OpAccessChain. First index must be 0.");
3771 const unsigned StartingIndex =
3772 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
3775 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
3776 Res.addUse(
I.getOperand(i).getReg());
3777 Res.constrainAllUses(
TII,
TRI, RBI);
3782bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
3784 unsigned Lim =
I.getNumExplicitOperands();
3785 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
3786 Register OpReg =
I.getOperand(i).getReg();
3787 MachineInstr *OpDefine =
MRI->getVRegDef(OpReg);
3790 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
3791 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
3798 MachineFunction *MF =
I.getMF();
3810 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3811 TII.get(SPIRV::OpSpecConstantOp))
3814 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
3816 GR.
add(OpDefine, MIB);
3822bool SPIRVInstructionSelector::selectDerivativeInst(
3823 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
3824 const unsigned DPdOpCode)
const {
3827 errorIfInstrOutsideShader(
I);
3832 Register SrcReg =
I.getOperand(2).getReg();
3837 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
3840 .
addUse(
I.getOperand(2).getReg());
3842 MachineIRBuilder MIRBuilder(
I);
3845 if (componentCount != 1)
3849 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
3850 Register ConvertToVReg =
MRI->createVirtualRegister(RegClass);
3851 Register DpdOpVReg =
MRI->createVirtualRegister(RegClass);
3853 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
3858 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
3863 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
3871bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
3872 SPIRVTypeInst ResType,
3873 MachineInstr &
I)
const {
3877 case Intrinsic::spv_load:
3878 return selectLoad(ResVReg, ResType,
I);
3879 case Intrinsic::spv_store:
3880 return selectStore(
I);
3881 case Intrinsic::spv_extractv:
3882 return selectExtractVal(ResVReg, ResType,
I);
3883 case Intrinsic::spv_insertv:
3884 return selectInsertVal(ResVReg, ResType,
I);
3885 case Intrinsic::spv_extractelt:
3886 return selectExtractElt(ResVReg, ResType,
I);
3887 case Intrinsic::spv_insertelt:
3888 return selectInsertElt(ResVReg, ResType,
I);
3889 case Intrinsic::spv_gep:
3890 return selectGEP(ResVReg, ResType,
I);
3891 case Intrinsic::spv_bitcast: {
3892 Register OpReg =
I.getOperand(2).getReg();
3893 SPIRVTypeInst OpType =
3897 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
3899 case Intrinsic::spv_unref_global:
3900 case Intrinsic::spv_init_global: {
3901 MachineInstr *
MI =
MRI->getVRegDef(
I.getOperand(1).getReg());
3902 MachineInstr *Init =
I.getNumExplicitOperands() > 2
3903 ?
MRI->getVRegDef(
I.getOperand(2).getReg())
3906 Register GVarVReg =
MI->getOperand(0).getReg();
3907 if (!selectGlobalValue(GVarVReg, *
MI, Init))
3912 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
3914 MI->eraseFromParent();
3918 case Intrinsic::spv_undef: {
3919 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3925 case Intrinsic::spv_const_composite: {
3927 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
3933 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
3935 MachineIRBuilder MIR(
I);
3937 MIR, SPIRV::OpConstantComposite, 3,
3938 SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,
3940 for (
auto *Instr : Instructions) {
3941 Instr->setDebugLoc(
I.getDebugLoc());
3946 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3953 case Intrinsic::spv_assign_name: {
3954 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
3955 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
3956 for (
unsigned i =
I.getNumExplicitDefs() + 2;
3957 i <
I.getNumExplicitOperands(); ++i) {
3958 MIB.
addImm(
I.getOperand(i).getImm());
3963 case Intrinsic::spv_switch: {
3964 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
3965 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
3966 if (
I.getOperand(i).isReg())
3967 MIB.
addReg(
I.getOperand(i).getReg());
3968 else if (
I.getOperand(i).isCImm())
3969 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
3970 else if (
I.getOperand(i).isMBB())
3971 MIB.
addMBB(
I.getOperand(i).getMBB());
3978 case Intrinsic::spv_loop_merge: {
3979 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
3980 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
3981 if (
I.getOperand(i).isMBB())
3982 MIB.
addMBB(
I.getOperand(i).getMBB());
3989 case Intrinsic::spv_loop_control_intel: {
3991 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
3992 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
3997 case Intrinsic::spv_selection_merge: {
3999 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4000 assert(
I.getOperand(1).isMBB() &&
4001 "operand 1 to spv_selection_merge must be a basic block");
4002 MIB.
addMBB(
I.getOperand(1).getMBB());
4003 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4007 case Intrinsic::spv_cmpxchg:
4008 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4009 case Intrinsic::spv_unreachable:
4010 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4013 case Intrinsic::spv_alloca:
4014 return selectFrameIndex(ResVReg, ResType,
I);
4015 case Intrinsic::spv_alloca_array:
4016 return selectAllocaArray(ResVReg, ResType,
I);
4017 case Intrinsic::spv_assume:
4019 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4020 .
addUse(
I.getOperand(1).getReg())
4025 case Intrinsic::spv_expect:
4027 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4030 .
addUse(
I.getOperand(2).getReg())
4031 .
addUse(
I.getOperand(3).getReg())
4036 case Intrinsic::arithmetic_fence:
4037 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4038 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4041 .
addUse(
I.getOperand(2).getReg())
4045 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4047 case Intrinsic::spv_thread_id:
4053 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4055 case Intrinsic::spv_thread_id_in_group:
4061 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4063 case Intrinsic::spv_group_id:
4069 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4071 case Intrinsic::spv_flattened_thread_id_in_group:
4078 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4080 case Intrinsic::spv_workgroup_size:
4081 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4083 case Intrinsic::spv_global_size:
4084 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4086 case Intrinsic::spv_global_offset:
4087 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4089 case Intrinsic::spv_num_workgroups:
4090 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4092 case Intrinsic::spv_subgroup_size:
4093 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4095 case Intrinsic::spv_num_subgroups:
4096 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4098 case Intrinsic::spv_subgroup_id:
4099 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4100 case Intrinsic::spv_subgroup_local_invocation_id:
4101 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4102 ResVReg, ResType,
I);
4103 case Intrinsic::spv_subgroup_max_size:
4104 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4106 case Intrinsic::spv_fdot:
4107 return selectFloatDot(ResVReg, ResType,
I);
4108 case Intrinsic::spv_udot:
4109 case Intrinsic::spv_sdot:
4110 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4112 return selectIntegerDot(ResVReg, ResType,
I,
4113 IID == Intrinsic::spv_sdot);
4114 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4115 case Intrinsic::spv_dot4add_i8packed:
4116 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4118 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4119 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4120 case Intrinsic::spv_dot4add_u8packed:
4121 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4123 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4124 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4125 case Intrinsic::spv_all:
4126 return selectAll(ResVReg, ResType,
I);
4127 case Intrinsic::spv_any:
4128 return selectAny(ResVReg, ResType,
I);
4129 case Intrinsic::spv_cross:
4130 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4131 case Intrinsic::spv_distance:
4132 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4133 case Intrinsic::spv_lerp:
4134 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4135 case Intrinsic::spv_length:
4136 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4137 case Intrinsic::spv_degrees:
4138 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4139 case Intrinsic::spv_faceforward:
4140 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4141 case Intrinsic::spv_frac:
4142 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4143 case Intrinsic::spv_isinf:
4144 return selectOpIsInf(ResVReg, ResType,
I);
4145 case Intrinsic::spv_isnan:
4146 return selectOpIsNan(ResVReg, ResType,
I);
4147 case Intrinsic::spv_normalize:
4148 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4149 case Intrinsic::spv_refract:
4150 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4151 case Intrinsic::spv_reflect:
4152 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4153 case Intrinsic::spv_rsqrt:
4154 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4155 case Intrinsic::spv_sign:
4156 return selectSign(ResVReg, ResType,
I);
4157 case Intrinsic::spv_smoothstep:
4158 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4159 case Intrinsic::spv_firstbituhigh:
4160 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4161 case Intrinsic::spv_firstbitshigh:
4162 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4163 case Intrinsic::spv_firstbitlow:
4164 return selectFirstBitLow(ResVReg, ResType,
I);
4165 case Intrinsic::spv_group_memory_barrier_with_group_sync: {
4167 buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent,
I);
4168 Register ScopeReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
4170 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpControlBarrier))
4177 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4178 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4179 SPIRV::StorageClass::StorageClass ResSC =
4183 "Generic storage class");
4184 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4192 case Intrinsic::spv_lifetime_start:
4193 case Intrinsic::spv_lifetime_end: {
4194 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
4195 : SPIRV::OpLifetimeStop;
4196 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
4197 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
4206 case Intrinsic::spv_saturate:
4207 return selectSaturate(ResVReg, ResType,
I);
4208 case Intrinsic::spv_nclamp:
4209 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
4210 case Intrinsic::spv_uclamp:
4211 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
4212 case Intrinsic::spv_sclamp:
4213 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
4214 case Intrinsic::spv_subgroup_prefix_bit_count:
4215 return selectWavePrefixBitCount(ResVReg, ResType,
I);
4216 case Intrinsic::spv_wave_active_countbits:
4217 return selectWaveActiveCountBits(ResVReg, ResType,
I);
4218 case Intrinsic::spv_wave_all_equal:
4219 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
4220 case Intrinsic::spv_wave_all:
4221 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
4222 case Intrinsic::spv_wave_any:
4223 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
4224 case Intrinsic::spv_subgroup_ballot:
4225 return selectWaveOpInst(ResVReg, ResType,
I,
4226 SPIRV::OpGroupNonUniformBallot);
4227 case Intrinsic::spv_wave_is_first_lane:
4228 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
4229 case Intrinsic::spv_wave_reduce_umax:
4230 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
4231 case Intrinsic::spv_wave_reduce_max:
4232 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
4233 case Intrinsic::spv_wave_reduce_umin:
4234 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
4235 case Intrinsic::spv_wave_reduce_min:
4236 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
4237 case Intrinsic::spv_wave_reduce_sum:
4238 return selectWaveReduceSum(ResVReg, ResType,
I);
4239 case Intrinsic::spv_wave_readlane:
4240 return selectWaveOpInst(ResVReg, ResType,
I,
4241 SPIRV::OpGroupNonUniformShuffle);
4242 case Intrinsic::spv_wave_prefix_sum:
4243 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
4244 case Intrinsic::spv_wave_prefix_product:
4245 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
4246 case Intrinsic::spv_step:
4247 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
4248 case Intrinsic::spv_radians:
4249 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
4253 case Intrinsic::instrprof_increment:
4254 case Intrinsic::instrprof_increment_step:
4255 case Intrinsic::instrprof_value_profile:
4258 case Intrinsic::spv_value_md:
4260 case Intrinsic::spv_resource_handlefrombinding: {
4261 return selectHandleFromBinding(ResVReg, ResType,
I);
4263 case Intrinsic::spv_resource_counterhandlefrombinding:
4264 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
4265 case Intrinsic::spv_resource_updatecounter:
4266 return selectUpdateCounter(ResVReg, ResType,
I);
4267 case Intrinsic::spv_resource_store_typedbuffer: {
4268 return selectImageWriteIntrinsic(
I);
4270 case Intrinsic::spv_resource_load_typedbuffer: {
4271 return selectReadImageIntrinsic(ResVReg, ResType,
I);
4273 case Intrinsic::spv_resource_sample:
4274 case Intrinsic::spv_resource_sample_clamp:
4275 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
4276 case Intrinsic::spv_resource_samplebias:
4277 case Intrinsic::spv_resource_samplebias_clamp:
4278 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
4279 case Intrinsic::spv_resource_samplegrad:
4280 case Intrinsic::spv_resource_samplegrad_clamp:
4281 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
4282 case Intrinsic::spv_resource_samplelevel:
4283 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
4284 case Intrinsic::spv_resource_samplecmp:
4285 case Intrinsic::spv_resource_samplecmp_clamp:
4286 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
4287 case Intrinsic::spv_resource_samplecmplevelzero:
4288 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
4289 case Intrinsic::spv_resource_gather:
4290 case Intrinsic::spv_resource_gather_cmp:
4291 return selectGatherIntrinsic(ResVReg, ResType,
I);
4292 case Intrinsic::spv_resource_getpointer: {
4293 return selectResourceGetPointer(ResVReg, ResType,
I);
4295 case Intrinsic::spv_pushconstant_getpointer: {
4296 return selectPushConstantGetPointer(ResVReg, ResType,
I);
4298 case Intrinsic::spv_discard: {
4299 return selectDiscard(ResVReg, ResType,
I);
4301 case Intrinsic::spv_resource_nonuniformindex: {
4302 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
4304 case Intrinsic::spv_unpackhalf2x16: {
4305 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
4307 case Intrinsic::spv_packhalf2x16: {
4308 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
4310 case Intrinsic::spv_ddx:
4311 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
4312 case Intrinsic::spv_ddy:
4313 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
4314 case Intrinsic::spv_ddx_coarse:
4315 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
4316 case Intrinsic::spv_ddy_coarse:
4317 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
4318 case Intrinsic::spv_ddx_fine:
4319 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
4320 case Intrinsic::spv_ddy_fine:
4321 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
4322 case Intrinsic::spv_fwidth:
4323 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
4325 std::string DiagMsg;
4326 raw_string_ostream OS(DiagMsg);
4328 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
4335bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
4336 SPIRVTypeInst ResType,
4337 MachineInstr &
I)
const {
4340 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
4347bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
4348 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4350 assert(Intr.getIntrinsicID() ==
4351 Intrinsic::spv_resource_counterhandlefrombinding);
4354 Register MainHandleReg = Intr.getOperand(2).getReg();
4356 assert(MainHandleDef->getIntrinsicID() ==
4357 Intrinsic::spv_resource_handlefrombinding);
4361 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(),
MRI);
4362 Register IndexReg = MainHandleDef->getOperand(5).getReg();
4363 std::string CounterName =
4368 MachineIRBuilder MIRBuilder(
I);
4370 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
4372 ArraySize, IndexReg, CounterName, MIRBuilder);
4374 return BuildCOPY(ResVReg, CounterVarReg,
I);
4377bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
4378 SPIRVTypeInst ResType,
4379 MachineInstr &
I)
const {
4381 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
4383 Register CounterHandleReg = Intr.getOperand(2).getReg();
4384 Register IncrReg = Intr.getOperand(3).getReg();
4391 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
4392 assert(CounterVarPointeeType &&
4393 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
4394 "Counter variable must be a struct");
4396 SPIRV::StorageClass::StorageBuffer &&
4397 "Counter variable must be in the storage buffer storage class");
4399 "Counter variable must have exactly 1 member in the struct");
4400 const SPIRVTypeInst MemberType =
4403 "Counter variable struct must have a single i32 member");
4407 MachineIRBuilder MIRBuilder(
I);
4409 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
4412 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
4418 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4421 .
addUse(CounterHandleReg)
4428 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
4431 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
4434 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
4443 return BuildCOPY(ResVReg, AtomicRes,
I);
4451 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
4459bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
4460 SPIRVTypeInst ResType,
4461 MachineInstr &
I)
const {
4469 Register ImageReg =
I.getOperand(2).getReg();
4471 Register NewImageReg =
MRI->createVirtualRegister(
MRI->getRegClass(ImageReg));
4477 Register IdxReg =
I.getOperand(3).getReg();
4479 MachineInstr &Pos =
I;
4481 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
4485bool SPIRVInstructionSelector::generateSampleImage(
4488 DebugLoc Loc, MachineInstr &Pos)
const {
4490 Register NewImageReg =
MRI->createVirtualRegister(
MRI->getRegClass(ImageReg));
4498 MRI->createVirtualRegister(
MRI->getRegClass(SamplerReg));
4499 if (!loadHandleBeforePosition(NewSamplerReg,
4505 MachineIRBuilder MIRBuilder(Pos);
4518 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
4519 ImOps.Lod.has_value();
4520 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
4521 : SPIRV::OpImageSampleImplicitLod;
4523 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
4524 : SPIRV::OpImageSampleDrefImplicitLod;
4533 MIB.
addUse(*ImOps.Compare);
4535 uint32_t ImageOperands = 0;
4537 ImageOperands |= SPIRV::ImageOperand::Bias;
4539 ImageOperands |= SPIRV::ImageOperand::Lod;
4540 if (ImOps.GradX && ImOps.GradY)
4541 ImageOperands |= SPIRV::ImageOperand::Grad;
4542 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
4544 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4547 "Non-constant offsets are not supported in sample instructions.");
4551 ImageOperands |= SPIRV::ImageOperand::MinLod;
4553 if (ImageOperands != 0) {
4554 MIB.
addImm(ImageOperands);
4555 if (ImageOperands & SPIRV::ImageOperand::Bias)
4557 if (ImageOperands & SPIRV::ImageOperand::Lod)
4559 if (ImageOperands & SPIRV::ImageOperand::Grad) {
4560 MIB.
addUse(*ImOps.GradX);
4561 MIB.
addUse(*ImOps.GradY);
4564 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4565 MIB.
addUse(*ImOps.Offset);
4566 if (ImageOperands & SPIRV::ImageOperand::MinLod)
4567 MIB.
addUse(*ImOps.MinLod);
4574bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
4575 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4576 Register ImageReg =
I.getOperand(2).getReg();
4577 Register SamplerReg =
I.getOperand(3).getReg();
4578 Register CoordinateReg =
I.getOperand(4).getReg();
4579 ImageOperands ImOps;
4580 if (
I.getNumOperands() > 5)
4581 ImOps.Offset =
I.getOperand(5).getReg();
4582 if (
I.getNumOperands() > 6)
4583 ImOps.MinLod =
I.getOperand(6).getReg();
4584 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4585 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4588bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
4589 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4590 Register ImageReg =
I.getOperand(2).getReg();
4591 Register SamplerReg =
I.getOperand(3).getReg();
4592 Register CoordinateReg =
I.getOperand(4).getReg();
4593 ImageOperands ImOps;
4594 ImOps.Bias =
I.getOperand(5).getReg();
4595 if (
I.getNumOperands() > 6)
4596 ImOps.Offset =
I.getOperand(6).getReg();
4597 if (
I.getNumOperands() > 7)
4598 ImOps.MinLod =
I.getOperand(7).getReg();
4599 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4600 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4603bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
4604 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4605 Register ImageReg =
I.getOperand(2).getReg();
4606 Register SamplerReg =
I.getOperand(3).getReg();
4607 Register CoordinateReg =
I.getOperand(4).getReg();
4608 ImageOperands ImOps;
4609 ImOps.GradX =
I.getOperand(5).getReg();
4610 ImOps.GradY =
I.getOperand(6).getReg();
4611 if (
I.getNumOperands() > 7)
4612 ImOps.Offset =
I.getOperand(7).getReg();
4613 if (
I.getNumOperands() > 8)
4614 ImOps.MinLod =
I.getOperand(8).getReg();
4615 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4616 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4619bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
4620 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4621 Register ImageReg =
I.getOperand(2).getReg();
4622 Register SamplerReg =
I.getOperand(3).getReg();
4623 Register CoordinateReg =
I.getOperand(4).getReg();
4624 ImageOperands ImOps;
4625 ImOps.Lod =
I.getOperand(5).getReg();
4626 if (
I.getNumOperands() > 6)
4627 ImOps.Offset =
I.getOperand(6).getReg();
4628 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4629 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4632bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
4633 SPIRVTypeInst ResType,
4634 MachineInstr &
I)
const {
4635 Register ImageReg =
I.getOperand(2).getReg();
4636 Register SamplerReg =
I.getOperand(3).getReg();
4637 Register CoordinateReg =
I.getOperand(4).getReg();
4638 ImageOperands ImOps;
4639 ImOps.Compare =
I.getOperand(5).getReg();
4640 if (
I.getNumOperands() > 6)
4641 ImOps.Offset =
I.getOperand(6).getReg();
4642 if (
I.getNumOperands() > 7)
4643 ImOps.MinLod =
I.getOperand(7).getReg();
4644 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4645 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4648bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
4649 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4650 Register ImageReg =
I.getOperand(2).getReg();
4651 Register SamplerReg =
I.getOperand(3).getReg();
4652 Register CoordinateReg =
I.getOperand(4).getReg();
4653 ImageOperands ImOps;
4654 ImOps.Compare =
I.getOperand(5).getReg();
4655 if (
I.getNumOperands() > 6)
4656 ImOps.Offset =
I.getOperand(6).getReg();
4659 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4660 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4663bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
4664 SPIRVTypeInst ResType,
4665 MachineInstr &
I)
const {
4666 Register ImageReg =
I.getOperand(2).getReg();
4667 Register SamplerReg =
I.getOperand(3).getReg();
4668 Register CoordinateReg =
I.getOperand(4).getReg();
4671 "ImageReg is not an image type.");
4676 ComponentOrCompareReg =
I.getOperand(5).getReg();
4677 OffsetReg =
I.getOperand(6).getReg();
4679 Register NewImageReg =
MRI->createVirtualRegister(
MRI->getRegClass(ImageReg));
4680 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
4684 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
4685 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
4686 Dim != SPIRV::Dim::DIM_Rect) {
4688 "Gather operations are only supported for 2D, Cube, and Rect images.");
4694 MRI->createVirtualRegister(
MRI->getRegClass(SamplerReg));
4695 if (!loadHandleBeforePosition(
4700 MachineIRBuilder MIRBuilder(
I);
4701 SPIRVTypeInst SampledImageType =
4706 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
4714 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
4716 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
4718 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4723 .
addUse(ComponentOrCompareReg);
4725 uint32_t ImageOperands = 0;
4726 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
4727 if (Dim == SPIRV::Dim::DIM_Cube) {
4729 "Gather operations with offset are not supported for Cube images.");
4733 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4735 ImageOperands |= SPIRV::ImageOperand::Offset;
4739 if (ImageOperands != 0) {
4740 MIB.
addImm(ImageOperands);
4742 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4750bool SPIRVInstructionSelector::generateImageReadOrFetch(
4755 "ImageReg is not an image type.");
4757 bool IsSignedInteger =
4762 bool IsFetch = (SampledOp.getImm() == 1);
4765 if (ResultSize == 4) {
4768 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
4774 if (IsSignedInteger)
4780 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
4784 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
4789 if (IsSignedInteger)
4793 if (ResultSize == 1) {
4802 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
4805bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
4806 SPIRVTypeInst ResType,
4807 MachineInstr &
I)
const {
4808 Register ResourcePtr =
I.getOperand(2).getReg();
4810 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
4819 MachineIRBuilder MIRBuilder(
I);
4821 Register IndexReg =
I.getOperand(3).getReg();
4824 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4834bool SPIRVInstructionSelector::selectPushConstantGetPointer(
4835 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4836 MRI->replaceRegWith(ResVReg,
I.getOperand(2).getReg());
4840bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
4841 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4842 Register ObjReg =
I.getOperand(2).getReg();
4843 if (!BuildCOPY(ResVReg, ObjReg,
I))
4853 decorateUsesAsNonUniform(ResVReg);
4857void SPIRVInstructionSelector::decorateUsesAsNonUniform(
4860 while (WorkList.
size() > 0) {
4864 bool IsDecorated =
false;
4865 for (MachineInstr &Use :
MRI->use_instructions(CurrentReg)) {
4866 if (
Use.getOpcode() == SPIRV::OpDecorate &&
4867 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
4873 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
4875 if (ResultReg == CurrentReg)
4883 SPIRV::Decoration::NonUniformEXT, {});
4888bool SPIRVInstructionSelector::extractSubvector(
4890 MachineInstr &InsertionPoint)
const {
4892 [[maybe_unused]] uint64_t InputSize =
4895 assert(InputSize > 1 &&
"The input must be a vector.");
4896 assert(ResultSize > 1 &&
"The result must be a vector.");
4897 assert(ResultSize < InputSize &&
4898 "Cannot extract more element than there are in the input.");
4901 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
4902 for (uint64_t
I = 0;
I < ResultSize;
I++) {
4903 Register ComponentReg =
MRI->createVirtualRegister(ScalarRegClass);
4905 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4914 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
4916 TII.get(SPIRV::OpCompositeConstruct))
4920 for (
Register ComponentReg : ComponentRegisters)
4921 MIB.
addUse(ComponentReg);
4926bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
4927 MachineInstr &
I)
const {
4934 Register ImageReg =
I.getOperand(1).getReg();
4936 Register NewImageReg =
MRI->createVirtualRegister(
MRI->getRegClass(ImageReg));
4942 Register CoordinateReg =
I.getOperand(2).getReg();
4943 Register DataReg =
I.getOperand(3).getReg();
4946 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
4954Register SPIRVInstructionSelector::buildPointerToResource(
4955 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
4956 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
4957 StringRef Name, MachineIRBuilder MIRBuilder)
const {
4959 if (ArraySize == 1) {
4960 SPIRVTypeInst PtrType =
4963 "SpirvResType did not have an explicit layout.");
4968 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
4969 SPIRVTypeInst VarPointerType =
4972 VarPointerType, Set,
Binding, Name, MIRBuilder);
4974 SPIRVTypeInst ResPointerType =
4987bool SPIRVInstructionSelector::selectFirstBitSet16(
4988 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4989 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
4991 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
4995 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
4998bool SPIRVInstructionSelector::selectFirstBitSet32(
5000 unsigned BitSetOpcode)
const {
5001 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5004 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5011bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
5013 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5020 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
5022 MachineIRBuilder MIRBuilder(
I);
5025 SPIRVTypeInst I64x2Type =
5027 SPIRVTypeInst Vec2ResType =
5030 std::vector<Register> PartialRegs;
5033 unsigned CurrentComponent = 0;
5034 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
5040 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5041 TII.get(SPIRV::OpVectorShuffle))
5046 .
addImm(CurrentComponent)
5047 .
addImm(CurrentComponent + 1);
5054 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType,
I, BitSetResult,
5055 BitSetOpcode, SwapPrimarySide))
5058 PartialRegs.push_back(SubVecBitSetReg);
5062 if (CurrentComponent != ComponentCount) {
5068 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
5069 SPIRV::OpVectorExtractDynamic))
5075 if (!selectFirstBitSet64(FinalElemBitSetReg,
BaseType,
I, FinalElemReg,
5076 BitSetOpcode, SwapPrimarySide))
5079 PartialRegs.push_back(FinalElemBitSetReg);
5084 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
5085 SPIRV::OpCompositeConstruct);
5088bool SPIRVInstructionSelector::selectFirstBitSet64(
5090 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5103 if (ComponentCount > 2) {
5104 return selectFirstBitSet64Overflow(ResVReg, ResType,
I, SrcReg,
5105 BitSetOpcode, SwapPrimarySide);
5109 MachineIRBuilder MIRBuilder(
I);
5111 BaseType, 2 * ComponentCount, MIRBuilder,
false);
5115 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
5121 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
5128 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
5131 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntZero},
5132 SPIRV::OpVectorExtractDynamic))
5134 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntOne},
5135 SPIRV::OpVectorExtractDynamic))
5139 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5140 TII.get(SPIRV::OpVectorShuffle))
5148 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
5154 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5155 TII.get(SPIRV::OpVectorShuffle))
5163 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
5183 SelectOp = SPIRV::OpSelectSISCond;
5184 AddOp = SPIRV::OpIAddS;
5192 SelectOp = SPIRV::OpSelectVIVCond;
5193 AddOp = SPIRV::OpIAddV;
5203 if (SwapPrimarySide) {
5204 PrimaryReg = LowReg;
5205 SecondaryReg = HighReg;
5206 PrimaryShiftReg = Reg0;
5207 SecondaryShiftReg = Reg32;
5212 if (!selectOpWithSrcs(BReg, BoolType,
I, {PrimaryReg, NegOneReg},
5218 if (!selectOpWithSrcs(TmpReg, ResType,
I, {
BReg, SecondaryReg, PrimaryReg},
5224 if (!selectOpWithSrcs(ValReg, ResType,
I,
5225 {
BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))
5228 return selectOpWithSrcs(ResVReg, ResType,
I, {ValReg, TmpReg}, AddOp);
5231bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
5232 SPIRVTypeInst ResType,
5234 bool IsSigned)
const {
5236 Register OpReg =
I.getOperand(2).getReg();
5239 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5240 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
5244 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5246 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5248 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5252 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
5256bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
5257 SPIRVTypeInst ResType,
5258 MachineInstr &
I)
const {
5260 Register OpReg =
I.getOperand(2).getReg();
5265 unsigned ExtendOpcode = SPIRV::OpUConvert;
5266 unsigned BitSetOpcode = GL::FindILsb;
5270 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5272 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5274 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5281bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
5282 SPIRVTypeInst ResType,
5283 MachineInstr &
I)
const {
5287 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
5290 .
addUse(
I.getOperand(2).getReg())
5293 unsigned Alignment =
I.getOperand(3).getImm();
5299bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
5300 SPIRVTypeInst ResType,
5301 MachineInstr &
I)
const {
5305 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
5308 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
5311 unsigned Alignment =
I.getOperand(2).getImm();
5318bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
5323 const MachineInstr *PrevI =
I.getPrevNode();
5325 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
5329 .
addMBB(
I.getOperand(0).getMBB())
5334 .
addMBB(
I.getOperand(0).getMBB())
5339bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
5350 const MachineInstr *NextI =
I.getNextNode();
5352 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
5358 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
5360 .
addUse(
I.getOperand(0).getReg())
5361 .
addMBB(
I.getOperand(1).getMBB())
5367bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
5368 MachineInstr &
I)
const {
5370 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
5372 const unsigned NumOps =
I.getNumOperands();
5373 for (
unsigned i = 1; i <
NumOps; i += 2) {
5374 MIB.
addUse(
I.getOperand(i + 0).getReg());
5375 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
5381bool SPIRVInstructionSelector::selectGlobalValue(
5382 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
5384 MachineIRBuilder MIRBuilder(
I);
5385 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
5388 std::string GlobalIdent;
5390 unsigned &
ID = UnnamedGlobalIDs[GV];
5392 ID = UnnamedGlobalIDs.
size();
5393 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
5419 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
5426 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
5429 MRI->setRegClass(FuncVReg, &SPIRV::pIDRegClass);
5431 MachineInstrBuilder MIB1 =
5432 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5435 MachineInstrBuilder MIB2 =
5437 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
5441 GR.
add(ConstVal, MIB2);
5449 MachineInstrBuilder MIB3 =
5450 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5453 GR.
add(ConstVal, MIB3);
5457 assert(NewReg != ResVReg);
5458 return BuildCOPY(ResVReg, NewReg,
I);
5468 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
5474 SPIRVTypeInst ResType =
5478 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
5483 if (
GlobalVar->isExternallyInitialized() &&
5484 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
5485 constexpr unsigned ReadWriteINTEL = 3u;
5488 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
5494bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
5495 SPIRVTypeInst ResType,
5496 MachineInstr &
I)
const {
5498 return selectExtInst(ResVReg, ResType,
I, CL::log10);
5506 MachineIRBuilder MIRBuilder(
I);
5511 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5514 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5516 .
add(
I.getOperand(1))
5521 ResType->
getOpcode() == SPIRV::OpTypeFloat);
5523 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
5531 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
5532 ? SPIRV::OpVectorTimesScalar
5543bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
5544 SPIRVTypeInst ResType,
5545 MachineInstr &
I)
const {
5561 MachineIRBuilder MIRBuilder(
I);
5564 ResType, MIRBuilder, SPIRV::StorageClass::Function);
5576 MachineBasicBlock &EntryBB =
I.getMF()->front();
5580 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
5583 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
5589 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5592 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
5595 .
add(
I.getOperand(
I.getNumExplicitDefs()))
5599 Register IntegralPartReg =
I.getOperand(1).getReg();
5600 if (IntegralPartReg.
isValid()) {
5602 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5613 assert(
false &&
"GLSL::Modf is deprecated.");
5624bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
5625 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5626 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5627 MachineIRBuilder MIRBuilder(
I);
5628 const SPIRVTypeInst Vec3Ty =
5631 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
5643 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5647 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
5648 Register LoadedRegister =
MRI->createVirtualRegister(&SPIRV::iIDRegClass);
5653 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5660 assert(
I.getOperand(2).isReg());
5661 const uint32_t ThreadId =
foldImm(
I.getOperand(2),
MRI);
5665 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5676bool SPIRVInstructionSelector::loadBuiltinInputID(
5677 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5678 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5679 MachineIRBuilder MIRBuilder(
I);
5681 ResType, MIRBuilder, SPIRV::StorageClass::Input);
5696 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5700 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5709SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
5710 MachineInstr &
I)
const {
5711 MachineIRBuilder MIRBuilder(
I);
5712 if (
Type->getOpcode() != SPIRV::OpTypeVector)
5716 if (VectorSize == 4)
5724bool SPIRVInstructionSelector::loadHandleBeforePosition(
5725 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
5726 MachineInstr &Pos)
const {
5729 Intrinsic::spv_resource_handlefrombinding);
5737 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
5738 MachineIRBuilder MIRBuilder(HandleDef);
5739 SPIRVTypeInst VarType = ResType;
5740 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
5742 if (IsStructuredBuffer) {
5748 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
5749 ArraySize, IndexReg, Name, MIRBuilder);
5753 uint32_t LoadOpcode =
5754 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
5764void SPIRVInstructionSelector::errorIfInstrOutsideShader(
5765 MachineInstr &
I)
const {
5767 std::string DiagMsg;
5768 raw_string_ostream OS(DiagMsg);
5769 I.print(OS,
true,
false,
false,
false);
5770 DiagMsg +=
" is only supported in shaders.\n";
5776InstructionSelector *
5780 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
unsigned const MachineRegisterInfo * MRI
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 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,...
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...
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
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.
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
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...