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) {
487 LLT RegType =
MRI.getType(
Reg);
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) {
798 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
799 Register SelectDstReg =
Def->getOperand(0).getReg();
803 Def->removeFromParent();
804 MRI->replaceRegWith(DstReg, SelectDstReg);
806 I.removeFromParent();
808 Res = selectImpl(
I, *CoverageInfo);
810 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
811 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
815 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
822 MRI->setRegClass(SrcReg,
MRI->getRegClass(DstReg));
823 MRI->replaceRegWith(SrcReg, DstReg);
827 }
else if (
I.getNumDefs() == 1) {
839 removeDeadInstruction(
I);
844 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
845 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
851 bool HasDefs =
I.getNumDefs() > 0;
854 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
855 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
856 if (spvSelect(ResVReg, ResType,
I)) {
858 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
869 case TargetOpcode::G_CONSTANT:
870 case TargetOpcode::G_FCONSTANT:
872 case TargetOpcode::G_SADDO:
873 case TargetOpcode::G_SSUBO:
880 MachineInstr &
I)
const {
881 const TargetRegisterClass *DstRC =
MRI->getRegClassOrNull(DestReg);
882 const TargetRegisterClass *SrcRC =
MRI->getRegClassOrNull(SrcReg);
883 if (DstRC != SrcRC && SrcRC)
884 MRI->setRegClass(DestReg, SrcRC);
885 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
892bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
893 SPIRVTypeInst ResType,
894 MachineInstr &
I)
const {
895 const unsigned Opcode =
I.getOpcode();
897 return selectImpl(
I, *CoverageInfo);
899 case TargetOpcode::G_CONSTANT:
900 case TargetOpcode::G_FCONSTANT:
901 return selectConst(ResVReg, ResType,
I);
902 case TargetOpcode::G_GLOBAL_VALUE:
903 return selectGlobalValue(ResVReg,
I);
904 case TargetOpcode::G_IMPLICIT_DEF:
905 return selectOpUndef(ResVReg, ResType,
I);
906 case TargetOpcode::G_FREEZE:
907 return selectFreeze(ResVReg, ResType,
I);
909 case TargetOpcode::G_INTRINSIC:
910 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
911 case TargetOpcode::G_INTRINSIC_CONVERGENT:
912 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
913 return selectIntrinsic(ResVReg, ResType,
I);
914 case TargetOpcode::G_BITREVERSE:
915 return selectBitreverse(ResVReg, ResType,
I);
917 case TargetOpcode::G_BUILD_VECTOR:
918 return selectBuildVector(ResVReg, ResType,
I);
919 case TargetOpcode::G_SPLAT_VECTOR:
920 return selectSplatVector(ResVReg, ResType,
I);
922 case TargetOpcode::G_SHUFFLE_VECTOR: {
923 MachineBasicBlock &BB = *
I.getParent();
924 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
927 .
addUse(
I.getOperand(1).getReg())
928 .
addUse(
I.getOperand(2).getReg());
929 for (
auto V :
I.getOperand(3).getShuffleMask())
934 case TargetOpcode::G_MEMMOVE:
935 case TargetOpcode::G_MEMCPY:
936 case TargetOpcode::G_MEMSET:
937 return selectMemOperation(ResVReg,
I);
939 case TargetOpcode::G_ICMP:
940 return selectICmp(ResVReg, ResType,
I);
941 case TargetOpcode::G_FCMP:
942 return selectFCmp(ResVReg, ResType,
I);
944 case TargetOpcode::G_FRAME_INDEX:
945 return selectFrameIndex(ResVReg, ResType,
I);
947 case TargetOpcode::G_LOAD:
948 return selectLoad(ResVReg, ResType,
I);
949 case TargetOpcode::G_STORE:
950 return selectStore(
I);
952 case TargetOpcode::G_BR:
953 return selectBranch(
I);
954 case TargetOpcode::G_BRCOND:
955 return selectBranchCond(
I);
957 case TargetOpcode::G_PHI:
958 return selectPhi(ResVReg,
I);
960 case TargetOpcode::G_FPTOSI:
961 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
962 case TargetOpcode::G_FPTOUI:
963 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
965 case TargetOpcode::G_FPTOSI_SAT:
966 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
967 case TargetOpcode::G_FPTOUI_SAT:
968 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
970 case TargetOpcode::G_SITOFP:
971 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
972 case TargetOpcode::G_UITOFP:
973 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
975 case TargetOpcode::G_CTPOP:
976 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitCount);
977 case TargetOpcode::G_SMIN:
978 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
979 case TargetOpcode::G_UMIN:
980 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
982 case TargetOpcode::G_SMAX:
983 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
984 case TargetOpcode::G_UMAX:
985 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
987 case TargetOpcode::G_SCMP:
988 return selectSUCmp(ResVReg, ResType,
I,
true);
989 case TargetOpcode::G_UCMP:
990 return selectSUCmp(ResVReg, ResType,
I,
false);
991 case TargetOpcode::G_LROUND:
992 case TargetOpcode::G_LLROUND: {
994 MRI->createVirtualRegister(
MRI->getRegClass(ResVReg),
"lround");
995 MRI->setRegClass(regForLround, &SPIRV::iIDRegClass);
997 regForLround, *(
I.getParent()->getParent()));
999 CL::round, GL::Round,
false);
1001 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1008 case TargetOpcode::G_STRICT_FMA:
1009 case TargetOpcode::G_FMA: {
1012 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1015 .
addUse(
I.getOperand(1).getReg())
1016 .
addUse(
I.getOperand(2).getReg())
1017 .
addUse(
I.getOperand(3).getReg())
1022 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1025 case TargetOpcode::G_STRICT_FLDEXP:
1026 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1028 case TargetOpcode::G_FPOW:
1029 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1030 case TargetOpcode::G_FPOWI:
1031 return selectExtInst(ResVReg, ResType,
I, CL::pown);
1033 case TargetOpcode::G_FEXP:
1034 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1035 case TargetOpcode::G_FEXP2:
1036 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1037 case TargetOpcode::G_FEXP10:
1038 return selectExp10(ResVReg, ResType,
I);
1040 case TargetOpcode::G_FMODF:
1041 return selectModf(ResVReg, ResType,
I);
1042 case TargetOpcode::G_FSINCOS:
1043 return selectSincos(ResVReg, ResType,
I);
1045 case TargetOpcode::G_FLOG:
1046 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1047 case TargetOpcode::G_FLOG2:
1048 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1049 case TargetOpcode::G_FLOG10:
1050 return selectLog10(ResVReg, ResType,
I);
1052 case TargetOpcode::G_FABS:
1053 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1054 case TargetOpcode::G_ABS:
1055 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1057 case TargetOpcode::G_FMINNUM:
1058 case TargetOpcode::G_FMINIMUM:
1059 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1060 case TargetOpcode::G_FMAXNUM:
1061 case TargetOpcode::G_FMAXIMUM:
1062 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1064 case TargetOpcode::G_FCOPYSIGN:
1065 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1067 case TargetOpcode::G_FCEIL:
1068 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1069 case TargetOpcode::G_FFLOOR:
1070 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1072 case TargetOpcode::G_FCOS:
1073 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1074 case TargetOpcode::G_FSIN:
1075 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1076 case TargetOpcode::G_FTAN:
1077 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1078 case TargetOpcode::G_FACOS:
1079 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1080 case TargetOpcode::G_FASIN:
1081 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1082 case TargetOpcode::G_FATAN:
1083 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1084 case TargetOpcode::G_FATAN2:
1085 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1086 case TargetOpcode::G_FCOSH:
1087 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1088 case TargetOpcode::G_FSINH:
1089 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1090 case TargetOpcode::G_FTANH:
1091 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1093 case TargetOpcode::G_STRICT_FSQRT:
1094 case TargetOpcode::G_FSQRT:
1095 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1097 case TargetOpcode::G_CTTZ:
1098 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1099 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1100 case TargetOpcode::G_CTLZ:
1101 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1102 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1104 case TargetOpcode::G_INTRINSIC_ROUND:
1105 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1106 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1107 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1108 case TargetOpcode::G_INTRINSIC_TRUNC:
1109 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1110 case TargetOpcode::G_FRINT:
1111 case TargetOpcode::G_FNEARBYINT:
1112 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1114 case TargetOpcode::G_SMULH:
1115 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1116 case TargetOpcode::G_UMULH:
1117 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1119 case TargetOpcode::G_SADDSAT:
1120 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1121 case TargetOpcode::G_UADDSAT:
1122 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1123 case TargetOpcode::G_SSUBSAT:
1124 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1125 case TargetOpcode::G_USUBSAT:
1126 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1128 case TargetOpcode::G_FFREXP:
1129 return selectFrexp(ResVReg, ResType,
I);
1131 case TargetOpcode::G_UADDO:
1132 return selectOverflowArith(ResVReg, ResType,
I,
1133 ResType->
getOpcode() == SPIRV::OpTypeVector
1134 ? SPIRV::OpIAddCarryV
1135 : SPIRV::OpIAddCarryS);
1136 case TargetOpcode::G_USUBO:
1137 return selectOverflowArith(ResVReg, ResType,
I,
1138 ResType->
getOpcode() == SPIRV::OpTypeVector
1139 ? SPIRV::OpISubBorrowV
1140 : SPIRV::OpISubBorrowS);
1141 case TargetOpcode::G_UMULO:
1142 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1143 case TargetOpcode::G_SMULO:
1144 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1146 case TargetOpcode::G_SEXT:
1147 return selectExt(ResVReg, ResType,
I,
true);
1148 case TargetOpcode::G_ANYEXT:
1149 case TargetOpcode::G_ZEXT:
1150 return selectExt(ResVReg, ResType,
I,
false);
1151 case TargetOpcode::G_TRUNC:
1152 return selectTrunc(ResVReg, ResType,
I);
1153 case TargetOpcode::G_FPTRUNC:
1154 case TargetOpcode::G_FPEXT:
1155 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1157 case TargetOpcode::G_PTRTOINT:
1158 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1159 case TargetOpcode::G_INTTOPTR:
1160 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1161 case TargetOpcode::G_BITCAST:
1162 return selectBitcast(ResVReg, ResType,
I);
1163 case TargetOpcode::G_ADDRSPACE_CAST:
1164 return selectAddrSpaceCast(ResVReg, ResType,
I);
1165 case TargetOpcode::G_PTR_ADD: {
1167 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1171 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1172 (*II).getOpcode() == TargetOpcode::COPY ||
1173 (*II).getOpcode() == SPIRV::OpVariable) &&
1176 bool IsGVInit =
false;
1178 UseIt =
MRI->use_instr_begin(
I.getOperand(0).getReg()),
1179 UseEnd =
MRI->use_instr_end();
1180 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1181 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1182 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1183 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1193 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1196 Register NewVReg =
MRI->createGenericVirtualRegister(
MRI->getType(GV));
1197 MRI->setRegClass(NewVReg,
MRI->getRegClass(GV));
1206 "incompatible result and operand types in a bitcast");
1208 MachineInstrBuilder MIB =
1209 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1216 : SPIRV::OpInBoundsPtrAccessChain))
1220 .
addUse(
I.getOperand(2).getReg())
1223 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1227 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1229 .
addUse(
I.getOperand(2).getReg())
1238 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1241 .
addImm(
static_cast<uint32_t
>(
1242 SPIRV::Opcode::InBoundsPtrAccessChain))
1245 .
addUse(
I.getOperand(2).getReg());
1250 case TargetOpcode::G_ATOMICRMW_OR:
1251 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1252 case TargetOpcode::G_ATOMICRMW_ADD:
1253 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1254 case TargetOpcode::G_ATOMICRMW_AND:
1255 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1256 case TargetOpcode::G_ATOMICRMW_MAX:
1257 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1258 case TargetOpcode::G_ATOMICRMW_MIN:
1259 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1260 case TargetOpcode::G_ATOMICRMW_SUB:
1261 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1262 case TargetOpcode::G_ATOMICRMW_XOR:
1263 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1264 case TargetOpcode::G_ATOMICRMW_UMAX:
1265 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1266 case TargetOpcode::G_ATOMICRMW_UMIN:
1267 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1268 case TargetOpcode::G_ATOMICRMW_XCHG:
1269 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1270 case TargetOpcode::G_ATOMIC_CMPXCHG:
1271 return selectAtomicCmpXchg(ResVReg, ResType,
I);
1273 case TargetOpcode::G_ATOMICRMW_FADD:
1274 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1275 case TargetOpcode::G_ATOMICRMW_FSUB:
1277 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1278 ResType->
getOpcode() == SPIRV::OpTypeVector
1280 : SPIRV::OpFNegate);
1281 case TargetOpcode::G_ATOMICRMW_FMIN:
1282 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1283 case TargetOpcode::G_ATOMICRMW_FMAX:
1284 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1286 case TargetOpcode::G_FENCE:
1287 return selectFence(
I);
1289 case TargetOpcode::G_STACKSAVE:
1290 return selectStackSave(ResVReg, ResType,
I);
1291 case TargetOpcode::G_STACKRESTORE:
1292 return selectStackRestore(
I);
1294 case TargetOpcode::G_UNMERGE_VALUES:
1300 case TargetOpcode::G_TRAP:
1301 case TargetOpcode::G_UBSANTRAP:
1302 case TargetOpcode::DBG_LABEL:
1304 case TargetOpcode::G_DEBUGTRAP:
1305 return selectDebugTrap(ResVReg, ResType,
I);
1312bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1313 SPIRVTypeInst ResType,
1314 MachineInstr &
I)
const {
1315 unsigned Opcode = SPIRV::OpNop;
1322bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1323 SPIRVTypeInst ResType,
1325 GL::GLSLExtInst GLInst,
1326 bool setMIFlags,
bool useMISrc,
1329 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1330 std::string DiagMsg;
1331 raw_string_ostream OS(DiagMsg);
1332 I.print(OS,
true,
false,
false,
false);
1333 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1336 return selectExtInst(ResVReg, ResType,
I,
1337 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1338 setMIFlags, useMISrc, SrcRegs);
1341bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1342 SPIRVTypeInst ResType,
1344 CL::OpenCLExtInst CLInst,
1345 bool setMIFlags,
bool useMISrc,
1347 return selectExtInst(ResVReg, ResType,
I,
1348 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1349 setMIFlags, useMISrc, SrcRegs);
1352bool SPIRVInstructionSelector::selectExtInst(
1353 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1354 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1356 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1357 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1358 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1362bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1363 SPIRVTypeInst ResType,
1366 bool setMIFlags,
bool useMISrc,
1369 for (
const auto &[InstructionSet, Opcode] : Insts) {
1373 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1376 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1381 const unsigned NumOps =
I.getNumOperands();
1384 I.getOperand(Index).getType() ==
1385 MachineOperand::MachineOperandType::MO_IntrinsicID)
1388 MIB.
add(
I.getOperand(Index));
1400bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1401 SPIRVTypeInst ResType,
1402 MachineInstr &
I)
const {
1403 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1404 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1405 for (
const auto &Ex : ExtInsts) {
1406 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1407 uint32_t Opcode = Ex.second;
1411 MachineIRBuilder MIRBuilder(
I);
1414 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1419 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1422 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1425 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1428 .
addImm(
static_cast<uint32_t
>(Ex.first))
1430 .
add(
I.getOperand(2))
1434 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1435 .
addDef(
I.getOperand(1).getReg())
1444bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1445 SPIRVTypeInst ResType,
1446 MachineInstr &
I)
const {
1447 Register CosResVReg =
I.getOperand(1).getReg();
1448 unsigned SrcIdx =
I.getNumExplicitDefs();
1453 MachineIRBuilder MIRBuilder(
I);
1455 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1460 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1463 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1465 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1468 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1470 .
add(
I.getOperand(SrcIdx))
1473 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1481 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1484 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1486 .
add(
I.getOperand(SrcIdx))
1488 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1491 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1493 .
add(
I.getOperand(SrcIdx))
1500bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1501 SPIRVTypeInst ResType,
1503 std::vector<Register> Srcs,
1504 unsigned Opcode)
const {
1505 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1515bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1516 SPIRVTypeInst ResType,
1518 unsigned Opcode)
const {
1520 Register SrcReg =
I.getOperand(1).getReg();
1523 MRI->def_instr_begin(SrcReg);
1524 DefIt !=
MRI->def_instr_end(); DefIt = std::next(DefIt)) {
1525 unsigned DefOpCode = DefIt->getOpcode();
1526 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1529 if (
auto *VRD =
getVRegDef(*
MRI, DefIt->getOperand(1).getReg()))
1530 DefOpCode = VRD->getOpcode();
1532 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1533 DefOpCode == TargetOpcode::G_CONSTANT ||
1534 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1540 uint32_t SpecOpcode = 0;
1542 case SPIRV::OpConvertPtrToU:
1543 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1545 case SPIRV::OpConvertUToPtr:
1546 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1551 TII.get(SPIRV::OpSpecConstantOp))
1561 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1565bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1566 SPIRVTypeInst ResType,
1567 MachineInstr &
I)
const {
1568 Register OpReg =
I.getOperand(1).getReg();
1569 SPIRVTypeInst OpType =
1573 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1583 if (
MemOp->isVolatile())
1584 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1585 if (
MemOp->isNonTemporal())
1586 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1588 if (!ST->isShader() &&
MemOp->getAlign().value())
1589 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1593 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1594 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1598 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1600 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1604 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1608 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1610 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1622 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1624 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1626 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1630bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1631 SPIRVTypeInst ResType,
1632 MachineInstr &
I)
const {
1634 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1639 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1640 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1642 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1644 MRI->createVirtualRegister(
MRI->getRegClass(HandleReg));
1646 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1650 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1651 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1652 I.getDebugLoc(),
I);
1656 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1660 if (!
I.getNumMemOperands()) {
1661 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1663 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1666 MachineIRBuilder MIRBuilder(
I);
1673bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1675 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1676 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1681 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1682 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1684 MRI->createVirtualRegister(
MRI->getRegClass(HandleReg));
1687 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1691 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1692 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1693 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1694 TII.get(SPIRV::OpImageWrite))
1700 if (sampledTypeIsSignedInteger(LLVMHandleType))
1703 BMI.constrainAllUses(
TII,
TRI, RBI);
1709 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpStore))
1712 if (!
I.getNumMemOperands()) {
1713 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1715 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1718 MachineIRBuilder MIRBuilder(
I);
1725bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
1726 SPIRVTypeInst ResType,
1727 MachineInstr &
I)
const {
1728 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1730 "llvm.stacksave intrinsic: this instruction requires the following "
1731 "SPIR-V extension: SPV_INTEL_variable_length_array",
1734 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
1741bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
1742 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1744 "llvm.stackrestore intrinsic: this instruction requires the following "
1745 "SPIR-V extension: SPV_INTEL_variable_length_array",
1747 if (!
I.getOperand(0).isReg())
1750 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
1751 .
addUse(
I.getOperand(0).getReg())
1757SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
1758 MachineIRBuilder MIRBuilder(
I);
1759 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1766 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
1770 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1771 Type *ArrTy = ArrayType::get(ValTy, Num);
1773 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
1776 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
1783 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
1786 .
addImm(SPIRV::StorageClass::UniformConstant)
1797bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
1800 Register DstReg =
I.getOperand(0).getReg();
1810 "Unable to determine pointee type size for OpCopyMemory");
1811 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
1812 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
1814 "OpCopyMemory requires the size to match the pointee type size");
1815 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
1818 if (
I.getNumMemOperands()) {
1819 MachineIRBuilder MIRBuilder(
I);
1826bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
1829 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
1830 .
addUse(
I.getOperand(0).getReg())
1832 .
addUse(
I.getOperand(2).getReg());
1833 if (
I.getNumMemOperands()) {
1834 MachineIRBuilder MIRBuilder(
I);
1841bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
1842 MachineInstr &
I)
const {
1843 Register SrcReg =
I.getOperand(1).getReg();
1844 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
1845 Register VarReg = getOrCreateMemSetGlobal(
I);
1848 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1850 ValTy,
I, SPIRV::StorageClass::UniformConstant);
1852 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
1856 if (!selectCopyMemory(
I, SrcReg))
1859 if (!selectCopyMemorySized(
I, SrcReg))
1862 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
1863 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
1868bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
1869 SPIRVTypeInst ResType,
1872 unsigned NegateOpcode)
const {
1874 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1877 Register ScopeReg = buildI32Constant(Scope,
I);
1879 Register Ptr =
I.getOperand(1).getReg();
1885 Register MemSemReg = buildI32Constant(MemSem ,
I);
1887 Register ValueReg =
I.getOperand(2).getReg();
1888 if (NegateOpcode != 0) {
1891 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
1896 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
1907bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
1908 unsigned ArgI =
I.getNumOperands() - 1;
1910 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
1911 SPIRVTypeInst SrcType =
1913 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
1915 "cannot select G_UNMERGE_VALUES with a non-vector argument");
1917 SPIRVTypeInst ScalarType =
1920 unsigned CurrentIndex = 0;
1921 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
1922 Register ResVReg =
I.getOperand(i).getReg();
1925 LLT ResLLT =
MRI->getType(ResVReg);
1931 ResType = ScalarType;
1937 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
1940 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1946 for (
unsigned j = 0;
j < NumElements; ++
j) {
1947 MIB.
addImm(CurrentIndex + j);
1949 CurrentIndex += NumElements;
1953 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
1965bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
1968 Register MemSemReg = buildI32Constant(MemSem,
I);
1970 uint32_t
Scope =
static_cast<uint32_t
>(
1972 Register ScopeReg = buildI32Constant(Scope,
I);
1974 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
1981bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
1982 SPIRVTypeInst ResType,
1984 unsigned Opcode)
const {
1985 Type *ResTy =
nullptr;
1989 "Not enough info to select the arithmetic with overflow instruction");
1992 "with overflow instruction");
1998 MachineIRBuilder MIRBuilder(
I);
2000 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2001 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2007 Register ZeroReg = buildZerosVal(ResType,
I);
2010 MRI->setRegClass(StructVReg, &SPIRV::IDRegClass);
2012 if (ResName.
size() > 0)
2017 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2020 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2021 MIB.
addUse(
I.getOperand(i).getReg());
2026 MRI->setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2027 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2029 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2030 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2037 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2038 .
addDef(
I.getOperand(1).getReg())
2046bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2047 SPIRVTypeInst ResType,
2048 MachineInstr &
I)
const {
2052 Register Ptr =
I.getOperand(2).getReg();
2055 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2058 ScopeReg = buildI32Constant(Scope,
I);
2060 unsigned ScSem =
static_cast<uint32_t
>(
2063 unsigned MemSemEq =
static_cast<uint32_t
>(
getMemSemantics(AO)) | ScSem;
2064 Register MemSemEqReg = buildI32Constant(MemSemEq,
I);
2066 unsigned MemSemNeq =
static_cast<uint32_t
>(
getMemSemantics(FO)) | ScSem;
2067 if (MemSemEq == MemSemNeq)
2068 MemSemNeqReg = MemSemEqReg;
2070 MemSemNeqReg = buildI32Constant(MemSemEq,
I);
2073 ScopeReg =
I.getOperand(5).getReg();
2074 MemSemEqReg =
I.getOperand(6).getReg();
2075 MemSemNeqReg =
I.getOperand(7).getReg();
2079 Register Val =
I.getOperand(4).getReg();
2083 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2102 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2109 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2121 case SPIRV::StorageClass::DeviceOnlyINTEL:
2122 case SPIRV::StorageClass::HostOnlyINTEL:
2131 bool IsGRef =
false;
2132 bool IsAllowedRefs =
2133 llvm::all_of(
MRI->use_instructions(ResVReg), [&IsGRef](
auto const &It) {
2134 unsigned Opcode = It.getOpcode();
2135 if (Opcode == SPIRV::OpConstantComposite ||
2136 Opcode == SPIRV::OpVariable ||
2137 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2138 return IsGRef = true;
2139 return Opcode == SPIRV::OpName;
2141 return IsAllowedRefs && IsGRef;
2144Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2145 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2147 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2151SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2153 uint32_t Opcode)
const {
2154 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2155 TII.get(SPIRV::OpSpecConstantOp))
2163SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2164 SPIRVTypeInst SrcPtrTy)
const {
2165 SPIRVTypeInst GenericPtrTy =
2167 Register Tmp =
MRI->createVirtualRegister(&SPIRV::pIDRegClass);
2169 SPIRV::StorageClass::Generic),
2171 MachineFunction *MF =
I.getParent()->getParent();
2173 MachineInstrBuilder MIB = buildSpecConstantOp(
2175 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2185bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2186 SPIRVTypeInst ResType,
2187 MachineInstr &
I)
const {
2191 Register SrcPtr =
I.getOperand(1).getReg();
2195 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2196 ResType->
getOpcode() != SPIRV::OpTypePointer)
2197 return BuildCOPY(ResVReg, SrcPtr,
I);
2207 unsigned SpecOpcode =
2209 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2212 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2219 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2221 .constrainAllUses(
TII,
TRI, RBI);
2223 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2225 buildSpecConstantOp(
2227 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2228 .constrainAllUses(
TII,
TRI, RBI);
2235 return BuildCOPY(ResVReg, SrcPtr,
I);
2237 if ((SrcSC == SPIRV::StorageClass::Function &&
2238 DstSC == SPIRV::StorageClass::Private) ||
2239 (DstSC == SPIRV::StorageClass::Function &&
2240 SrcSC == SPIRV::StorageClass::Private))
2241 return BuildCOPY(ResVReg, SrcPtr,
I);
2245 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2248 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2251 SPIRVTypeInst GenericPtrTy =
2270 return selectUnOp(ResVReg, ResType,
I,
2271 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2273 return selectUnOp(ResVReg, ResType,
I,
2274 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2276 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2278 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2288 return SPIRV::OpFOrdEqual;
2290 return SPIRV::OpFOrdGreaterThanEqual;
2292 return SPIRV::OpFOrdGreaterThan;
2294 return SPIRV::OpFOrdLessThanEqual;
2296 return SPIRV::OpFOrdLessThan;
2298 return SPIRV::OpFOrdNotEqual;
2300 return SPIRV::OpOrdered;
2302 return SPIRV::OpFUnordEqual;
2304 return SPIRV::OpFUnordGreaterThanEqual;
2306 return SPIRV::OpFUnordGreaterThan;
2308 return SPIRV::OpFUnordLessThanEqual;
2310 return SPIRV::OpFUnordLessThan;
2312 return SPIRV::OpFUnordNotEqual;
2314 return SPIRV::OpUnordered;
2324 return SPIRV::OpIEqual;
2326 return SPIRV::OpINotEqual;
2328 return SPIRV::OpSGreaterThanEqual;
2330 return SPIRV::OpSGreaterThan;
2332 return SPIRV::OpSLessThanEqual;
2334 return SPIRV::OpSLessThan;
2336 return SPIRV::OpUGreaterThanEqual;
2338 return SPIRV::OpUGreaterThan;
2340 return SPIRV::OpULessThanEqual;
2342 return SPIRV::OpULessThan;
2351 return SPIRV::OpPtrEqual;
2353 return SPIRV::OpPtrNotEqual;
2364 return SPIRV::OpLogicalEqual;
2366 return SPIRV::OpLogicalNotEqual;
2400bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2401 SPIRVTypeInst ResType,
2403 unsigned OpAnyOrAll)
const {
2404 assert(
I.getNumOperands() == 3);
2405 assert(
I.getOperand(2).isReg());
2407 Register InputRegister =
I.getOperand(2).getReg();
2414 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2415 if (IsBoolTy && !IsVectorTy) {
2416 assert(ResVReg ==
I.getOperand(0).getReg());
2417 return BuildCOPY(ResVReg, InputRegister,
I);
2421 unsigned SpirvNotEqualId =
2422 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2424 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2429 IsBoolTy ? InputRegister
2437 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2439 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2456bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2457 SPIRVTypeInst ResType,
2458 MachineInstr &
I)
const {
2459 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2462bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2463 SPIRVTypeInst ResType,
2464 MachineInstr &
I)
const {
2465 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2469bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2470 SPIRVTypeInst ResType,
2471 MachineInstr &
I)
const {
2472 assert(
I.getNumOperands() == 4);
2473 assert(
I.getOperand(2).isReg());
2474 assert(
I.getOperand(3).isReg());
2476 [[maybe_unused]] SPIRVTypeInst VecType =
2481 "dot product requires a vector of at least 2 components");
2483 [[maybe_unused]] SPIRVTypeInst EltType =
2492 .
addUse(
I.getOperand(2).getReg())
2493 .
addUse(
I.getOperand(3).getReg())
2498bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2499 SPIRVTypeInst ResType,
2502 assert(
I.getNumOperands() == 4);
2503 assert(
I.getOperand(2).isReg());
2504 assert(
I.getOperand(3).isReg());
2507 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2511 .
addUse(
I.getOperand(2).getReg())
2512 .
addUse(
I.getOperand(3).getReg())
2519bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2520 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2521 assert(
I.getNumOperands() == 4);
2522 assert(
I.getOperand(2).isReg());
2523 assert(
I.getOperand(3).isReg());
2527 Register Vec0 =
I.getOperand(2).getReg();
2528 Register Vec1 =
I.getOperand(3).getReg();
2532 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2541 "dot product requires a vector of at least 2 components");
2544 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2554 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2565 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2577bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2578 SPIRVTypeInst ResType,
2579 MachineInstr &
I)
const {
2581 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2584 .
addUse(
I.getOperand(2).getReg())
2589bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2590 SPIRVTypeInst ResType,
2591 MachineInstr &
I)
const {
2593 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2596 .
addUse(
I.getOperand(2).getReg())
2601template <
bool Signed>
2602bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2603 SPIRVTypeInst ResType,
2604 MachineInstr &
I)
const {
2605 assert(
I.getNumOperands() == 5);
2606 assert(
I.getOperand(2).isReg());
2607 assert(
I.getOperand(3).isReg());
2608 assert(
I.getOperand(4).isReg());
2611 Register Acc =
I.getOperand(2).getReg();
2615 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2617 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
2622 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
2625 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2637template <
bool Signed>
2638bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
2639 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2640 assert(
I.getNumOperands() == 5);
2641 assert(
I.getOperand(2).isReg());
2642 assert(
I.getOperand(3).isReg());
2643 assert(
I.getOperand(4).isReg());
2646 Register Acc =
I.getOperand(2).getReg();
2652 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
2656 for (
unsigned i = 0; i < 4; i++) {
2658 Register AElt =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2668 Register BElt =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2679 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
2687 Register MaskMul =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2698 i < 3 ?
MRI->createVirtualRegister(&SPIRV::IDRegClass) : ResVReg;
2699 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2714bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
2715 SPIRVTypeInst ResType,
2716 MachineInstr &
I)
const {
2717 assert(
I.getNumOperands() == 3);
2718 assert(
I.getOperand(2).isReg());
2720 Register VZero = buildZerosValF(ResType,
I);
2721 Register VOne = buildOnesValF(ResType,
I);
2723 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
2726 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2728 .
addUse(
I.getOperand(2).getReg())
2735bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
2736 SPIRVTypeInst ResType,
2737 MachineInstr &
I)
const {
2738 assert(
I.getNumOperands() == 3);
2739 assert(
I.getOperand(2).isReg());
2741 Register InputRegister =
I.getOperand(2).getReg();
2743 auto &
DL =
I.getDebugLoc();
2753 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
2755 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
2757 ?
MRI->createVirtualRegister(&SPIRV::IDRegClass)
2763 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2768 if (NeedsConversion) {
2769 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
2780bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
2781 SPIRVTypeInst ResType,
2783 unsigned Opcode)
const {
2787 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2793 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
2794 BMI.addUse(
I.getOperand(J).getReg());
2801bool SPIRVInstructionSelector::selectWaveActiveCountBits(
2802 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2807 if (!selectWaveOpInst(BallotReg, BallotType,
I,
2808 SPIRV::OpGroupNonUniformBallot))
2813 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2818 .
addImm(SPIRV::GroupOperation::Reduce)
2827 if (
Type->getOpcode() != SPIRV::OpTypeVector)
2831 return Type->getOperand(2).getImm();
2834bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
2835 SPIRVTypeInst ResType,
2836 MachineInstr &
I)
const {
2841 Register InputReg =
I.getOperand(2).getReg();
2846 bool IsVector = NumElems > 1;
2849 SPIRVTypeInst ElemInputType = InputType;
2850 SPIRVTypeInst ElemBoolType = ResType;
2863 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
2864 SPIRV::OpGroupNonUniformAllEqual);
2869 ElementResults.
reserve(NumElems);
2871 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
2884 ElemInput = Extracted;
2890 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
2901 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
2912bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
2913 SPIRVTypeInst ResType,
2914 MachineInstr &
I)
const {
2916 assert(
I.getNumOperands() == 3);
2918 auto Op =
I.getOperand(2);
2930 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
2941 Register BallotVReg =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2952 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2956 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
2963bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
2964 SPIRVTypeInst ResType,
2966 bool IsUnsigned)
const {
2967 return selectWaveReduce(
2968 ResVReg, ResType,
I, IsUnsigned,
2969 [&](
Register InputRegister,
bool IsUnsigned) {
2970 const bool IsFloatTy =
2972 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
2973 : SPIRV::OpGroupNonUniformSMax;
2974 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
2978bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
2979 SPIRVTypeInst ResType,
2981 bool IsUnsigned)
const {
2982 return selectWaveReduce(
2983 ResVReg, ResType,
I, IsUnsigned,
2984 [&](
Register InputRegister,
bool IsUnsigned) {
2985 const bool IsFloatTy =
2987 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
2988 : SPIRV::OpGroupNonUniformSMin;
2989 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
2993bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
2994 SPIRVTypeInst ResType,
2995 MachineInstr &
I)
const {
2996 return selectWaveReduce(ResVReg, ResType,
I,
false,
2997 [&](
Register InputRegister,
bool IsUnsigned) {
2999 InputRegister, SPIRV::OpTypeFloat);
3000 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3001 : SPIRV::OpGroupNonUniformIAdd;
3005template <
typename PickOpcodeFn>
3006bool SPIRVInstructionSelector::selectWaveReduce(
3007 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3008 PickOpcodeFn &&PickOpcode)
const {
3009 assert(
I.getNumOperands() == 3);
3010 assert(
I.getOperand(2).isReg());
3012 Register InputRegister =
I.getOperand(2).getReg();
3019 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3025 .
addImm(SPIRV::GroupOperation::Reduce)
3026 .
addUse(
I.getOperand(2).getReg())
3031bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3032 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3033 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3034 [&](
Register InputRegister,
bool IsUnsigned) {
3036 InputRegister, SPIRV::OpTypeFloat);
3038 ? SPIRV::OpGroupNonUniformFAdd
3039 : SPIRV::OpGroupNonUniformIAdd;
3043bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3044 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3045 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3046 [&](
Register InputRegister,
bool IsUnsigned) {
3048 InputRegister, SPIRV::OpTypeFloat);
3050 ? SPIRV::OpGroupNonUniformFMul
3051 : SPIRV::OpGroupNonUniformIMul;
3055template <
typename PickOpcodeFn>
3056bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3057 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3058 PickOpcodeFn &&PickOpcode)
const {
3059 assert(
I.getNumOperands() == 3);
3060 assert(
I.getOperand(2).isReg());
3062 Register InputRegister =
I.getOperand(2).getReg();
3069 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3075 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3076 .
addUse(
I.getOperand(2).getReg())
3081bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3082 SPIRVTypeInst ResType,
3083 MachineInstr &
I)
const {
3085 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3088 .
addUse(
I.getOperand(1).getReg())
3093bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3094 SPIRVTypeInst ResType,
3095 MachineInstr &
I)
const {
3101 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3103 Register OpReg =
I.getOperand(1).getReg();
3104 if (MachineInstr *Def =
MRI->getVRegDef(OpReg)) {
3105 if (
Def->getOpcode() == TargetOpcode::COPY)
3106 Def =
MRI->getVRegDef(
Def->getOperand(1).getReg());
3108 switch (
Def->getOpcode()) {
3109 case SPIRV::ASSIGN_TYPE:
3110 if (MachineInstr *AssignToDef =
3111 MRI->getVRegDef(
Def->getOperand(1).getReg())) {
3112 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3113 Reg =
Def->getOperand(2).getReg();
3116 case SPIRV::OpUndef:
3117 Reg =
Def->getOperand(1).getReg();
3120 unsigned DestOpCode;
3122 DestOpCode = SPIRV::OpConstantNull;
3124 DestOpCode = TargetOpcode::COPY;
3127 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3128 .
addDef(
I.getOperand(0).getReg())
3136bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3137 SPIRVTypeInst ResType,
3138 MachineInstr &
I)
const {
3140 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3142 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3146 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3151 for (
unsigned i =
I.getNumExplicitDefs();
3152 i <
I.getNumExplicitOperands() && IsConst; ++i)
3156 if (!IsConst &&
N < 2)
3158 "There must be at least two constituent operands in a vector");
3161 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3162 TII.get(IsConst ? SPIRV::OpConstantComposite
3163 : SPIRV::OpCompositeConstruct))
3166 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3167 MIB.
addUse(
I.getOperand(i).getReg());
3172bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3173 SPIRVTypeInst ResType,
3174 MachineInstr &
I)
const {
3176 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3178 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3184 if (!
I.getOperand(
OpIdx).isReg())
3191 if (!IsConst &&
N < 2)
3193 "There must be at least two constituent operands in a vector");
3196 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3197 TII.get(IsConst ? SPIRV::OpConstantComposite
3198 : SPIRV::OpCompositeConstruct))
3201 for (
unsigned i = 0; i <
N; ++i)
3207bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3208 SPIRVTypeInst ResType,
3209 MachineInstr &
I)
const {
3214 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3216 Opcode = SPIRV::OpDemoteToHelperInvocation;
3218 Opcode = SPIRV::OpKill;
3220 if (MachineInstr *NextI =
I.getNextNode()) {
3222 NextI->eraseFromParent();
3232bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3233 SPIRVTypeInst ResType,
unsigned CmpOpc,
3234 MachineInstr &
I)
const {
3235 Register Cmp0 =
I.getOperand(2).getReg();
3236 Register Cmp1 =
I.getOperand(3).getReg();
3239 "CMP operands should have the same type");
3240 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3250bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3251 SPIRVTypeInst ResType,
3252 MachineInstr &
I)
const {
3253 auto Pred =
I.getOperand(1).getPredicate();
3256 Register CmpOperand =
I.getOperand(2).getReg();
3263 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3267SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3268 SPIRVTypeInst ResType)
const {
3270 SPIRVTypeInst SpvI32Ty =
3273 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3280 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3283 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3286 .
addImm(APInt(32, Val).getZExtValue());
3288 GR.
add(ConstInt,
MI);
3293bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
3294 SPIRVTypeInst ResType,
3295 MachineInstr &
I)
const {
3297 return selectCmp(ResVReg, ResType, CmpOp,
I);
3300bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
3301 SPIRVTypeInst ResType,
3302 MachineInstr &
I)
const {
3304 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
3311 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
3312 ResType->
getOpcode() != SPIRV::OpTypeFloat)
3315 MachineIRBuilder MIRBuilder(
I);
3317 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
3323 "only float operands supported by GLSL extended math");
3326 MIRBuilder, SpirvScalarType);
3328 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
3329 ? SPIRV::OpVectorTimesScalar
3332 if (!selectOpWithSrcs(ArgReg, ResType,
I,
3333 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
3335 if (!selectExtInst(ResVReg, ResType,
I,
3336 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
3346Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
3347 MachineInstr &
I)
const {
3350 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3355bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
3361 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
3369 if (
Def->getOpcode() == SPIRV::OpConstantNull)
3372 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
3373 Def->getOpcode() == SPIRV::OpConstantI)
3382 MachineInstr *
Def =
MRI->getVRegDef(
Reg);
3386 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
3387 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
3389 Intrinsic::spv_const_composite)) {
3390 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
3391 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
3392 if (!IsZero(
Def->getOperand(i).getReg()))
3401Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
3402 MachineInstr &
I)
const {
3406 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3411Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
3412 MachineInstr &
I)
const {
3416 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3422 SPIRVTypeInst ResType,
3423 MachineInstr &
I)
const {
3427 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3432bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
3433 SPIRVTypeInst ResType,
3434 MachineInstr &
I)
const {
3435 Register SelectFirstArg =
I.getOperand(2).getReg();
3436 Register SelectSecondArg =
I.getOperand(3).getReg();
3445 SPIRV::OpTypeVector;
3452 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
3453 }
else if (IsPtrTy) {
3454 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
3456 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
3460 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
3461 }
else if (IsPtrTy) {
3462 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
3464 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3467 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3470 .
addUse(
I.getOperand(1).getReg())
3479bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
3480 SPIRVTypeInst ResType,
3482 MachineInstr &InsertAt,
3483 bool IsSigned)
const {
3485 Register ZeroReg = buildZerosVal(ResType, InsertAt);
3486 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
3487 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
3489 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3501bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
3502 SPIRVTypeInst ResType,
3503 MachineInstr &
I,
bool IsSigned,
3504 unsigned Opcode)
const {
3505 Register SrcReg =
I.getOperand(1).getReg();
3511 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
3516 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
3518 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
3521bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
3522 SPIRVTypeInst ResType, MachineInstr &
I,
3523 bool IsSigned)
const {
3524 Register SrcReg =
I.getOperand(1).getReg();
3526 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
3530 if (ResType == SrcType)
3531 return BuildCOPY(ResVReg, SrcReg,
I);
3533 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3534 return selectUnOp(ResVReg, ResType,
I, Opcode);
3537bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
3538 SPIRVTypeInst ResType,
3540 bool IsSigned)
const {
3541 MachineIRBuilder MIRBuilder(
I);
3542 MachineRegisterInfo *
MRI = MIRBuilder.getMRI();
3557 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
3560 .
addUse(
I.getOperand(1).getReg())
3561 .
addUse(
I.getOperand(2).getReg())
3567 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
3570 .
addUse(
I.getOperand(1).getReg())
3571 .
addUse(
I.getOperand(2).getReg())
3579 unsigned SelectOpcode =
3580 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
3585 .
addUse(buildOnesVal(
true, ResType,
I))
3586 .
addUse(buildZerosVal(ResType,
I))
3593 .
addUse(buildOnesVal(
false, ResType,
I))
3598bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
3601 SPIRVTypeInst IntTy,
3602 SPIRVTypeInst BoolTy)
const {
3605 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
3606 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
3608 Register One = buildOnesVal(
false, IntTy,
I);
3616 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
3625bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
3626 SPIRVTypeInst ResType,
3627 MachineInstr &
I)
const {
3628 Register IntReg =
I.getOperand(1).getReg();
3631 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
3632 if (ArgType == ResType)
3633 return BuildCOPY(ResVReg, IntReg,
I);
3635 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3636 return selectUnOp(ResVReg, ResType,
I, Opcode);
3639bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
3640 SPIRVTypeInst ResType,
3641 MachineInstr &
I)
const {
3642 unsigned Opcode =
I.getOpcode();
3643 unsigned TpOpcode = ResType->
getOpcode();
3645 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
3646 assert(Opcode == TargetOpcode::G_CONSTANT &&
3647 I.getOperand(1).getCImm()->isZero());
3648 MachineBasicBlock &DepMBB =
I.getMF()->front();
3651 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
3658 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
3661bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
3662 SPIRVTypeInst ResType,
3663 MachineInstr &
I)
const {
3664 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3671bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
3672 SPIRVTypeInst ResType,
3673 MachineInstr &
I)
const {
3675 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
3679 .
addUse(
I.getOperand(3).getReg())
3681 .
addUse(
I.getOperand(2).getReg());
3682 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
3688bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
3689 SPIRVTypeInst ResType,
3690 MachineInstr &
I)
const {
3691 Type *MaybeResTy =
nullptr;
3696 "Expected aggregate type for extractv instruction");
3698 SPIRV::AccessQualifier::ReadWrite,
false);
3702 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3705 .
addUse(
I.getOperand(2).getReg());
3706 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
3712bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
3713 SPIRVTypeInst ResType,
3714 MachineInstr &
I)
const {
3716 return selectInsertVal(ResVReg, ResType,
I);
3718 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
3721 .
addUse(
I.getOperand(2).getReg())
3722 .
addUse(
I.getOperand(3).getReg())
3723 .
addUse(
I.getOperand(4).getReg())
3728bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
3729 SPIRVTypeInst ResType,
3730 MachineInstr &
I)
const {
3732 return selectExtractVal(ResVReg, ResType,
I);
3734 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
3737 .
addUse(
I.getOperand(2).getReg())
3738 .
addUse(
I.getOperand(3).getReg())
3743bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
3744 SPIRVTypeInst ResType,
3745 MachineInstr &
I)
const {
3746 const bool IsGEPInBounds =
I.getOperand(2).getImm();
3752 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
3753 : SPIRV::OpAccessChain)
3754 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
3755 :
SPIRV::OpPtrAccessChain);
3757 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3761 .
addUse(
I.getOperand(3).getReg());
3763 (Opcode == SPIRV::OpPtrAccessChain ||
3764 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
3766 "Cannot translate GEP to OpAccessChain. First index must be 0.");
3769 const unsigned StartingIndex =
3770 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
3773 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
3774 Res.addUse(
I.getOperand(i).getReg());
3775 Res.constrainAllUses(
TII,
TRI, RBI);
3780bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
3782 unsigned Lim =
I.getNumExplicitOperands();
3783 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
3784 Register OpReg =
I.getOperand(i).getReg();
3785 MachineInstr *OpDefine =
MRI->getVRegDef(OpReg);
3788 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
3789 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
3796 MachineFunction *MF =
I.getMF();
3808 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3809 TII.get(SPIRV::OpSpecConstantOp))
3812 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
3814 GR.
add(OpDefine, MIB);
3820bool SPIRVInstructionSelector::selectDerivativeInst(
3821 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
3822 const unsigned DPdOpCode)
const {
3825 errorIfInstrOutsideShader(
I);
3830 Register SrcReg =
I.getOperand(2).getReg();
3835 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
3838 .
addUse(
I.getOperand(2).getReg());
3840 MachineIRBuilder MIRBuilder(
I);
3843 if (componentCount != 1)
3847 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
3848 Register ConvertToVReg =
MRI->createVirtualRegister(RegClass);
3849 Register DpdOpVReg =
MRI->createVirtualRegister(RegClass);
3851 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
3856 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
3861 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
3869bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
3870 SPIRVTypeInst ResType,
3871 MachineInstr &
I)
const {
3875 case Intrinsic::spv_load:
3876 return selectLoad(ResVReg, ResType,
I);
3877 case Intrinsic::spv_store:
3878 return selectStore(
I);
3879 case Intrinsic::spv_extractv:
3880 return selectExtractVal(ResVReg, ResType,
I);
3881 case Intrinsic::spv_insertv:
3882 return selectInsertVal(ResVReg, ResType,
I);
3883 case Intrinsic::spv_extractelt:
3884 return selectExtractElt(ResVReg, ResType,
I);
3885 case Intrinsic::spv_insertelt:
3886 return selectInsertElt(ResVReg, ResType,
I);
3887 case Intrinsic::spv_gep:
3888 return selectGEP(ResVReg, ResType,
I);
3889 case Intrinsic::spv_bitcast: {
3890 Register OpReg =
I.getOperand(2).getReg();
3891 SPIRVTypeInst OpType =
3895 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
3897 case Intrinsic::spv_unref_global:
3898 case Intrinsic::spv_init_global: {
3899 MachineInstr *
MI =
MRI->getVRegDef(
I.getOperand(1).getReg());
3900 MachineInstr *Init =
I.getNumExplicitOperands() > 2
3901 ?
MRI->getVRegDef(
I.getOperand(2).getReg())
3904 Register GVarVReg =
MI->getOperand(0).getReg();
3905 if (!selectGlobalValue(GVarVReg, *
MI, Init))
3910 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
3912 MI->eraseFromParent();
3916 case Intrinsic::spv_undef: {
3917 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3923 case Intrinsic::spv_const_composite: {
3925 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
3931 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
3933 MachineIRBuilder MIR(
I);
3935 MIR, SPIRV::OpConstantComposite, 3,
3936 SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,
3938 for (
auto *Instr : Instructions) {
3939 Instr->setDebugLoc(
I.getDebugLoc());
3944 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3951 case Intrinsic::spv_assign_name: {
3952 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
3953 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
3954 for (
unsigned i =
I.getNumExplicitDefs() + 2;
3955 i <
I.getNumExplicitOperands(); ++i) {
3956 MIB.
addImm(
I.getOperand(i).getImm());
3961 case Intrinsic::spv_switch: {
3962 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
3963 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
3964 if (
I.getOperand(i).isReg())
3965 MIB.
addReg(
I.getOperand(i).getReg());
3966 else if (
I.getOperand(i).isCImm())
3967 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
3968 else if (
I.getOperand(i).isMBB())
3969 MIB.
addMBB(
I.getOperand(i).getMBB());
3976 case Intrinsic::spv_loop_merge: {
3977 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
3978 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
3979 if (
I.getOperand(i).isMBB())
3980 MIB.
addMBB(
I.getOperand(i).getMBB());
3987 case Intrinsic::spv_loop_control_intel: {
3989 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
3990 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
3995 case Intrinsic::spv_selection_merge: {
3997 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
3998 assert(
I.getOperand(1).isMBB() &&
3999 "operand 1 to spv_selection_merge must be a basic block");
4000 MIB.
addMBB(
I.getOperand(1).getMBB());
4001 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4005 case Intrinsic::spv_cmpxchg:
4006 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4007 case Intrinsic::spv_unreachable:
4008 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4011 case Intrinsic::spv_alloca:
4012 return selectFrameIndex(ResVReg, ResType,
I);
4013 case Intrinsic::spv_alloca_array:
4014 return selectAllocaArray(ResVReg, ResType,
I);
4015 case Intrinsic::spv_assume:
4017 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4018 .
addUse(
I.getOperand(1).getReg())
4023 case Intrinsic::spv_expect:
4025 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4028 .
addUse(
I.getOperand(2).getReg())
4029 .
addUse(
I.getOperand(3).getReg())
4034 case Intrinsic::arithmetic_fence:
4035 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4036 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4039 .
addUse(
I.getOperand(2).getReg())
4043 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4045 case Intrinsic::spv_thread_id:
4051 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4053 case Intrinsic::spv_thread_id_in_group:
4059 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4061 case Intrinsic::spv_group_id:
4067 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4069 case Intrinsic::spv_flattened_thread_id_in_group:
4076 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4078 case Intrinsic::spv_workgroup_size:
4079 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4081 case Intrinsic::spv_global_size:
4082 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4084 case Intrinsic::spv_global_offset:
4085 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4087 case Intrinsic::spv_num_workgroups:
4088 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4090 case Intrinsic::spv_subgroup_size:
4091 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4093 case Intrinsic::spv_num_subgroups:
4094 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4096 case Intrinsic::spv_subgroup_id:
4097 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4098 case Intrinsic::spv_subgroup_local_invocation_id:
4099 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4100 ResVReg, ResType,
I);
4101 case Intrinsic::spv_subgroup_max_size:
4102 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4104 case Intrinsic::spv_fdot:
4105 return selectFloatDot(ResVReg, ResType,
I);
4106 case Intrinsic::spv_udot:
4107 case Intrinsic::spv_sdot:
4108 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4110 return selectIntegerDot(ResVReg, ResType,
I,
4111 IID == Intrinsic::spv_sdot);
4112 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4113 case Intrinsic::spv_dot4add_i8packed:
4114 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4116 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4117 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4118 case Intrinsic::spv_dot4add_u8packed:
4119 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4121 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4122 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4123 case Intrinsic::spv_all:
4124 return selectAll(ResVReg, ResType,
I);
4125 case Intrinsic::spv_any:
4126 return selectAny(ResVReg, ResType,
I);
4127 case Intrinsic::spv_cross:
4128 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4129 case Intrinsic::spv_distance:
4130 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4131 case Intrinsic::spv_lerp:
4132 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4133 case Intrinsic::spv_length:
4134 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4135 case Intrinsic::spv_degrees:
4136 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4137 case Intrinsic::spv_faceforward:
4138 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4139 case Intrinsic::spv_frac:
4140 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4141 case Intrinsic::spv_isinf:
4142 return selectOpIsInf(ResVReg, ResType,
I);
4143 case Intrinsic::spv_isnan:
4144 return selectOpIsNan(ResVReg, ResType,
I);
4145 case Intrinsic::spv_normalize:
4146 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4147 case Intrinsic::spv_refract:
4148 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4149 case Intrinsic::spv_reflect:
4150 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4151 case Intrinsic::spv_rsqrt:
4152 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4153 case Intrinsic::spv_sign:
4154 return selectSign(ResVReg, ResType,
I);
4155 case Intrinsic::spv_smoothstep:
4156 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4157 case Intrinsic::spv_firstbituhigh:
4158 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4159 case Intrinsic::spv_firstbitshigh:
4160 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4161 case Intrinsic::spv_firstbitlow:
4162 return selectFirstBitLow(ResVReg, ResType,
I);
4163 case Intrinsic::spv_group_memory_barrier_with_group_sync: {
4165 buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent,
I);
4166 Register ScopeReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
4168 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpControlBarrier))
4175 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4176 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4177 SPIRV::StorageClass::StorageClass ResSC =
4181 "Generic storage class");
4182 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4190 case Intrinsic::spv_lifetime_start:
4191 case Intrinsic::spv_lifetime_end: {
4192 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
4193 : SPIRV::OpLifetimeStop;
4194 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
4195 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
4204 case Intrinsic::spv_saturate:
4205 return selectSaturate(ResVReg, ResType,
I);
4206 case Intrinsic::spv_nclamp:
4207 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
4208 case Intrinsic::spv_uclamp:
4209 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
4210 case Intrinsic::spv_sclamp:
4211 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
4212 case Intrinsic::spv_subgroup_prefix_bit_count:
4213 return selectWavePrefixBitCount(ResVReg, ResType,
I);
4214 case Intrinsic::spv_wave_active_countbits:
4215 return selectWaveActiveCountBits(ResVReg, ResType,
I);
4216 case Intrinsic::spv_wave_all_equal:
4217 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
4218 case Intrinsic::spv_wave_all:
4219 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
4220 case Intrinsic::spv_wave_any:
4221 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
4222 case Intrinsic::spv_subgroup_ballot:
4223 return selectWaveOpInst(ResVReg, ResType,
I,
4224 SPIRV::OpGroupNonUniformBallot);
4225 case Intrinsic::spv_wave_is_first_lane:
4226 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
4227 case Intrinsic::spv_wave_reduce_umax:
4228 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
4229 case Intrinsic::spv_wave_reduce_max:
4230 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
4231 case Intrinsic::spv_wave_reduce_umin:
4232 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
4233 case Intrinsic::spv_wave_reduce_min:
4234 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
4235 case Intrinsic::spv_wave_reduce_sum:
4236 return selectWaveReduceSum(ResVReg, ResType,
I);
4237 case Intrinsic::spv_wave_readlane:
4238 return selectWaveOpInst(ResVReg, ResType,
I,
4239 SPIRV::OpGroupNonUniformShuffle);
4240 case Intrinsic::spv_wave_prefix_sum:
4241 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
4242 case Intrinsic::spv_wave_prefix_product:
4243 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
4244 case Intrinsic::spv_step:
4245 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
4246 case Intrinsic::spv_radians:
4247 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
4251 case Intrinsic::instrprof_increment:
4252 case Intrinsic::instrprof_increment_step:
4253 case Intrinsic::instrprof_value_profile:
4256 case Intrinsic::spv_value_md:
4258 case Intrinsic::spv_resource_handlefrombinding: {
4259 return selectHandleFromBinding(ResVReg, ResType,
I);
4261 case Intrinsic::spv_resource_counterhandlefrombinding:
4262 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
4263 case Intrinsic::spv_resource_updatecounter:
4264 return selectUpdateCounter(ResVReg, ResType,
I);
4265 case Intrinsic::spv_resource_store_typedbuffer: {
4266 return selectImageWriteIntrinsic(
I);
4268 case Intrinsic::spv_resource_load_typedbuffer: {
4269 return selectReadImageIntrinsic(ResVReg, ResType,
I);
4271 case Intrinsic::spv_resource_sample:
4272 case Intrinsic::spv_resource_sample_clamp:
4273 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
4274 case Intrinsic::spv_resource_samplebias:
4275 case Intrinsic::spv_resource_samplebias_clamp:
4276 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
4277 case Intrinsic::spv_resource_samplegrad:
4278 case Intrinsic::spv_resource_samplegrad_clamp:
4279 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
4280 case Intrinsic::spv_resource_samplelevel:
4281 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
4282 case Intrinsic::spv_resource_samplecmp:
4283 case Intrinsic::spv_resource_samplecmp_clamp:
4284 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
4285 case Intrinsic::spv_resource_samplecmplevelzero:
4286 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
4287 case Intrinsic::spv_resource_gather:
4288 case Intrinsic::spv_resource_gather_cmp:
4289 return selectGatherIntrinsic(ResVReg, ResType,
I);
4290 case Intrinsic::spv_resource_getpointer: {
4291 return selectResourceGetPointer(ResVReg, ResType,
I);
4293 case Intrinsic::spv_pushconstant_getpointer: {
4294 return selectPushConstantGetPointer(ResVReg, ResType,
I);
4296 case Intrinsic::spv_discard: {
4297 return selectDiscard(ResVReg, ResType,
I);
4299 case Intrinsic::spv_resource_nonuniformindex: {
4300 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
4302 case Intrinsic::spv_unpackhalf2x16: {
4303 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
4305 case Intrinsic::spv_packhalf2x16: {
4306 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
4308 case Intrinsic::spv_ddx:
4309 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
4310 case Intrinsic::spv_ddy:
4311 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
4312 case Intrinsic::spv_ddx_coarse:
4313 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
4314 case Intrinsic::spv_ddy_coarse:
4315 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
4316 case Intrinsic::spv_ddx_fine:
4317 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
4318 case Intrinsic::spv_ddy_fine:
4319 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
4320 case Intrinsic::spv_fwidth:
4321 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
4323 std::string DiagMsg;
4324 raw_string_ostream OS(DiagMsg);
4326 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
4333bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
4334 SPIRVTypeInst ResType,
4335 MachineInstr &
I)
const {
4338 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
4345bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
4346 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4348 assert(Intr.getIntrinsicID() ==
4349 Intrinsic::spv_resource_counterhandlefrombinding);
4352 Register MainHandleReg = Intr.getOperand(2).getReg();
4354 assert(MainHandleDef->getIntrinsicID() ==
4355 Intrinsic::spv_resource_handlefrombinding);
4359 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(),
MRI);
4360 Register IndexReg = MainHandleDef->getOperand(5).getReg();
4361 std::string CounterName =
4366 MachineIRBuilder MIRBuilder(
I);
4368 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
4370 ArraySize, IndexReg, CounterName, MIRBuilder);
4372 return BuildCOPY(ResVReg, CounterVarReg,
I);
4375bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
4376 SPIRVTypeInst ResType,
4377 MachineInstr &
I)
const {
4379 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
4381 Register CounterHandleReg = Intr.getOperand(2).getReg();
4382 Register IncrReg = Intr.getOperand(3).getReg();
4389 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
4390 assert(CounterVarPointeeType &&
4391 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
4392 "Counter variable must be a struct");
4394 SPIRV::StorageClass::StorageBuffer &&
4395 "Counter variable must be in the storage buffer storage class");
4397 "Counter variable must have exactly 1 member in the struct");
4398 const SPIRVTypeInst MemberType =
4401 "Counter variable struct must have a single i32 member");
4405 MachineIRBuilder MIRBuilder(
I);
4407 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
4410 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
4416 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4419 .
addUse(CounterHandleReg)
4426 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
4429 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
4432 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
4441 return BuildCOPY(ResVReg, AtomicRes,
I);
4449 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
4457bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
4458 SPIRVTypeInst ResType,
4459 MachineInstr &
I)
const {
4467 Register ImageReg =
I.getOperand(2).getReg();
4469 Register NewImageReg =
MRI->createVirtualRegister(
MRI->getRegClass(ImageReg));
4475 Register IdxReg =
I.getOperand(3).getReg();
4477 MachineInstr &Pos =
I;
4479 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
4483bool SPIRVInstructionSelector::generateSampleImage(
4486 DebugLoc Loc, MachineInstr &Pos)
const {
4488 Register NewImageReg =
MRI->createVirtualRegister(
MRI->getRegClass(ImageReg));
4496 MRI->createVirtualRegister(
MRI->getRegClass(SamplerReg));
4497 if (!loadHandleBeforePosition(NewSamplerReg,
4503 MachineIRBuilder MIRBuilder(Pos);
4516 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
4517 ImOps.Lod.has_value();
4518 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
4519 : SPIRV::OpImageSampleImplicitLod;
4521 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
4522 : SPIRV::OpImageSampleDrefImplicitLod;
4531 MIB.
addUse(*ImOps.Compare);
4533 uint32_t ImageOperands = 0;
4535 ImageOperands |= SPIRV::ImageOperand::Bias;
4537 ImageOperands |= SPIRV::ImageOperand::Lod;
4538 if (ImOps.GradX && ImOps.GradY)
4539 ImageOperands |= SPIRV::ImageOperand::Grad;
4540 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
4542 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4545 "Non-constant offsets are not supported in sample instructions.");
4549 ImageOperands |= SPIRV::ImageOperand::MinLod;
4551 if (ImageOperands != 0) {
4552 MIB.
addImm(ImageOperands);
4553 if (ImageOperands & SPIRV::ImageOperand::Bias)
4555 if (ImageOperands & SPIRV::ImageOperand::Lod)
4557 if (ImageOperands & SPIRV::ImageOperand::Grad) {
4558 MIB.
addUse(*ImOps.GradX);
4559 MIB.
addUse(*ImOps.GradY);
4562 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4563 MIB.
addUse(*ImOps.Offset);
4564 if (ImageOperands & SPIRV::ImageOperand::MinLod)
4565 MIB.
addUse(*ImOps.MinLod);
4572bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
4573 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4574 Register ImageReg =
I.getOperand(2).getReg();
4575 Register SamplerReg =
I.getOperand(3).getReg();
4576 Register CoordinateReg =
I.getOperand(4).getReg();
4577 ImageOperands ImOps;
4578 if (
I.getNumOperands() > 5)
4579 ImOps.Offset =
I.getOperand(5).getReg();
4580 if (
I.getNumOperands() > 6)
4581 ImOps.MinLod =
I.getOperand(6).getReg();
4582 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4583 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4586bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
4587 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4588 Register ImageReg =
I.getOperand(2).getReg();
4589 Register SamplerReg =
I.getOperand(3).getReg();
4590 Register CoordinateReg =
I.getOperand(4).getReg();
4591 ImageOperands ImOps;
4592 ImOps.Bias =
I.getOperand(5).getReg();
4593 if (
I.getNumOperands() > 6)
4594 ImOps.Offset =
I.getOperand(6).getReg();
4595 if (
I.getNumOperands() > 7)
4596 ImOps.MinLod =
I.getOperand(7).getReg();
4597 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4598 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4601bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
4602 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4603 Register ImageReg =
I.getOperand(2).getReg();
4604 Register SamplerReg =
I.getOperand(3).getReg();
4605 Register CoordinateReg =
I.getOperand(4).getReg();
4606 ImageOperands ImOps;
4607 ImOps.GradX =
I.getOperand(5).getReg();
4608 ImOps.GradY =
I.getOperand(6).getReg();
4609 if (
I.getNumOperands() > 7)
4610 ImOps.Offset =
I.getOperand(7).getReg();
4611 if (
I.getNumOperands() > 8)
4612 ImOps.MinLod =
I.getOperand(8).getReg();
4613 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4614 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4617bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
4618 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4619 Register ImageReg =
I.getOperand(2).getReg();
4620 Register SamplerReg =
I.getOperand(3).getReg();
4621 Register CoordinateReg =
I.getOperand(4).getReg();
4622 ImageOperands ImOps;
4623 ImOps.Lod =
I.getOperand(5).getReg();
4624 if (
I.getNumOperands() > 6)
4625 ImOps.Offset =
I.getOperand(6).getReg();
4626 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4627 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4630bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
4631 SPIRVTypeInst ResType,
4632 MachineInstr &
I)
const {
4633 Register ImageReg =
I.getOperand(2).getReg();
4634 Register SamplerReg =
I.getOperand(3).getReg();
4635 Register CoordinateReg =
I.getOperand(4).getReg();
4636 ImageOperands ImOps;
4637 ImOps.Compare =
I.getOperand(5).getReg();
4638 if (
I.getNumOperands() > 6)
4639 ImOps.Offset =
I.getOperand(6).getReg();
4640 if (
I.getNumOperands() > 7)
4641 ImOps.MinLod =
I.getOperand(7).getReg();
4642 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4643 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4646bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
4647 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4648 Register ImageReg =
I.getOperand(2).getReg();
4649 Register SamplerReg =
I.getOperand(3).getReg();
4650 Register CoordinateReg =
I.getOperand(4).getReg();
4651 ImageOperands ImOps;
4652 ImOps.Compare =
I.getOperand(5).getReg();
4653 if (
I.getNumOperands() > 6)
4654 ImOps.Offset =
I.getOperand(6).getReg();
4657 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4658 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4661bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
4662 SPIRVTypeInst ResType,
4663 MachineInstr &
I)
const {
4664 Register ImageReg =
I.getOperand(2).getReg();
4665 Register SamplerReg =
I.getOperand(3).getReg();
4666 Register CoordinateReg =
I.getOperand(4).getReg();
4669 "ImageReg is not an image type.");
4674 ComponentOrCompareReg =
I.getOperand(5).getReg();
4675 OffsetReg =
I.getOperand(6).getReg();
4677 Register NewImageReg =
MRI->createVirtualRegister(
MRI->getRegClass(ImageReg));
4678 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
4682 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
4683 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
4684 Dim != SPIRV::Dim::DIM_Rect) {
4686 "Gather operations are only supported for 2D, Cube, and Rect images.");
4692 MRI->createVirtualRegister(
MRI->getRegClass(SamplerReg));
4693 if (!loadHandleBeforePosition(
4698 MachineIRBuilder MIRBuilder(
I);
4699 SPIRVTypeInst SampledImageType =
4704 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
4712 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
4714 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
4716 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4721 .
addUse(ComponentOrCompareReg);
4723 uint32_t ImageOperands = 0;
4724 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
4725 if (Dim == SPIRV::Dim::DIM_Cube) {
4727 "Gather operations with offset are not supported for Cube images.");
4731 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4733 ImageOperands |= SPIRV::ImageOperand::Offset;
4737 if (ImageOperands != 0) {
4738 MIB.
addImm(ImageOperands);
4740 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4748bool SPIRVInstructionSelector::generateImageReadOrFetch(
4753 "ImageReg is not an image type.");
4755 bool IsSignedInteger =
4760 bool IsFetch = (SampledOp.getImm() == 1);
4763 if (ResultSize == 4) {
4766 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
4772 if (IsSignedInteger)
4778 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
4782 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
4787 if (IsSignedInteger)
4791 if (ResultSize == 1) {
4800 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
4803bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
4804 SPIRVTypeInst ResType,
4805 MachineInstr &
I)
const {
4806 Register ResourcePtr =
I.getOperand(2).getReg();
4808 if (RegType->
getOpcode() == SPIRV::OpTypeImage) {
4817 MachineIRBuilder MIRBuilder(
I);
4819 Register IndexReg =
I.getOperand(3).getReg();
4822 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4832bool SPIRVInstructionSelector::selectPushConstantGetPointer(
4833 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4834 MRI->replaceRegWith(ResVReg,
I.getOperand(2).getReg());
4838bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
4839 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4840 Register ObjReg =
I.getOperand(2).getReg();
4841 if (!BuildCOPY(ResVReg, ObjReg,
I))
4851 decorateUsesAsNonUniform(ResVReg);
4855void SPIRVInstructionSelector::decorateUsesAsNonUniform(
4858 while (WorkList.
size() > 0) {
4862 bool IsDecorated =
false;
4863 for (MachineInstr &Use :
MRI->use_instructions(CurrentReg)) {
4864 if (
Use.getOpcode() == SPIRV::OpDecorate &&
4865 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
4871 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
4873 if (ResultReg == CurrentReg)
4881 SPIRV::Decoration::NonUniformEXT, {});
4886bool SPIRVInstructionSelector::extractSubvector(
4888 MachineInstr &InsertionPoint)
const {
4890 [[maybe_unused]] uint64_t InputSize =
4893 assert(InputSize > 1 &&
"The input must be a vector.");
4894 assert(ResultSize > 1 &&
"The result must be a vector.");
4895 assert(ResultSize < InputSize &&
4896 "Cannot extract more element than there are in the input.");
4899 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
4900 for (uint64_t
I = 0;
I < ResultSize;
I++) {
4901 Register ComponentReg =
MRI->createVirtualRegister(ScalarRegClass);
4903 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4912 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
4914 TII.get(SPIRV::OpCompositeConstruct))
4918 for (
Register ComponentReg : ComponentRegisters)
4919 MIB.
addUse(ComponentReg);
4924bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
4925 MachineInstr &
I)
const {
4932 Register ImageReg =
I.getOperand(1).getReg();
4934 Register NewImageReg =
MRI->createVirtualRegister(
MRI->getRegClass(ImageReg));
4940 Register CoordinateReg =
I.getOperand(2).getReg();
4941 Register DataReg =
I.getOperand(3).getReg();
4944 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
4952Register SPIRVInstructionSelector::buildPointerToResource(
4953 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
4954 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
4955 StringRef Name, MachineIRBuilder MIRBuilder)
const {
4957 if (ArraySize == 1) {
4958 SPIRVTypeInst PtrType =
4961 "SpirvResType did not have an explicit layout.");
4966 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
4967 SPIRVTypeInst VarPointerType =
4970 VarPointerType, Set,
Binding, Name, MIRBuilder);
4972 SPIRVTypeInst ResPointerType =
4985bool SPIRVInstructionSelector::selectFirstBitSet16(
4986 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4987 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
4989 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
4993 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
4996bool SPIRVInstructionSelector::selectFirstBitSet32(
4998 unsigned BitSetOpcode)
const {
4999 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5002 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5009bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
5011 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5018 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
5020 MachineIRBuilder MIRBuilder(
I);
5023 SPIRVTypeInst I64x2Type =
5025 SPIRVTypeInst Vec2ResType =
5028 std::vector<Register> PartialRegs;
5031 unsigned CurrentComponent = 0;
5032 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
5038 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5039 TII.get(SPIRV::OpVectorShuffle))
5044 .
addImm(CurrentComponent)
5045 .
addImm(CurrentComponent + 1);
5052 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType,
I, BitSetResult,
5053 BitSetOpcode, SwapPrimarySide))
5056 PartialRegs.push_back(SubVecBitSetReg);
5060 if (CurrentComponent != ComponentCount) {
5066 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
5067 SPIRV::OpVectorExtractDynamic))
5073 if (!selectFirstBitSet64(FinalElemBitSetReg,
BaseType,
I, FinalElemReg,
5074 BitSetOpcode, SwapPrimarySide))
5077 PartialRegs.push_back(FinalElemBitSetReg);
5082 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
5083 SPIRV::OpCompositeConstruct);
5086bool SPIRVInstructionSelector::selectFirstBitSet64(
5088 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5101 if (ComponentCount > 2) {
5102 return selectFirstBitSet64Overflow(ResVReg, ResType,
I, SrcReg,
5103 BitSetOpcode, SwapPrimarySide);
5107 MachineIRBuilder MIRBuilder(
I);
5109 BaseType, 2 * ComponentCount, MIRBuilder,
false);
5113 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
5119 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
5126 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
5129 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntZero},
5130 SPIRV::OpVectorExtractDynamic))
5132 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntOne},
5133 SPIRV::OpVectorExtractDynamic))
5137 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5138 TII.get(SPIRV::OpVectorShuffle))
5146 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
5152 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5153 TII.get(SPIRV::OpVectorShuffle))
5161 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
5181 SelectOp = SPIRV::OpSelectSISCond;
5182 AddOp = SPIRV::OpIAddS;
5190 SelectOp = SPIRV::OpSelectVIVCond;
5191 AddOp = SPIRV::OpIAddV;
5201 if (SwapPrimarySide) {
5202 PrimaryReg = LowReg;
5203 SecondaryReg = HighReg;
5204 PrimaryShiftReg = Reg0;
5205 SecondaryShiftReg = Reg32;
5210 if (!selectOpWithSrcs(BReg, BoolType,
I, {PrimaryReg, NegOneReg},
5216 if (!selectOpWithSrcs(TmpReg, ResType,
I, {
BReg, SecondaryReg, PrimaryReg},
5222 if (!selectOpWithSrcs(ValReg, ResType,
I,
5223 {
BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))
5226 return selectOpWithSrcs(ResVReg, ResType,
I, {ValReg, TmpReg}, AddOp);
5229bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
5230 SPIRVTypeInst ResType,
5232 bool IsSigned)
const {
5234 Register OpReg =
I.getOperand(2).getReg();
5237 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5238 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
5242 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5244 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5246 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5250 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
5254bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
5255 SPIRVTypeInst ResType,
5256 MachineInstr &
I)
const {
5258 Register OpReg =
I.getOperand(2).getReg();
5263 unsigned ExtendOpcode = SPIRV::OpUConvert;
5264 unsigned BitSetOpcode = GL::FindILsb;
5268 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5270 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5272 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5279bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
5280 SPIRVTypeInst ResType,
5281 MachineInstr &
I)
const {
5285 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
5288 .
addUse(
I.getOperand(2).getReg())
5291 unsigned Alignment =
I.getOperand(3).getImm();
5297bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
5298 SPIRVTypeInst ResType,
5299 MachineInstr &
I)
const {
5303 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
5306 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
5309 unsigned Alignment =
I.getOperand(2).getImm();
5316bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
5321 const MachineInstr *PrevI =
I.getPrevNode();
5323 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
5327 .
addMBB(
I.getOperand(0).getMBB())
5332 .
addMBB(
I.getOperand(0).getMBB())
5337bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
5348 const MachineInstr *NextI =
I.getNextNode();
5350 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
5356 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
5358 .
addUse(
I.getOperand(0).getReg())
5359 .
addMBB(
I.getOperand(1).getMBB())
5365bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
5366 MachineInstr &
I)
const {
5368 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
5370 const unsigned NumOps =
I.getNumOperands();
5371 for (
unsigned i = 1; i <
NumOps; i += 2) {
5372 MIB.
addUse(
I.getOperand(i + 0).getReg());
5373 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
5379bool SPIRVInstructionSelector::selectGlobalValue(
5380 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
5382 MachineIRBuilder MIRBuilder(
I);
5383 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
5386 std::string GlobalIdent;
5388 unsigned &
ID = UnnamedGlobalIDs[GV];
5390 ID = UnnamedGlobalIDs.
size();
5391 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
5417 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
5424 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
5427 MRI->setRegClass(FuncVReg, &SPIRV::pIDRegClass);
5429 MachineInstrBuilder MIB1 =
5430 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5433 MachineInstrBuilder MIB2 =
5435 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
5439 GR.
add(ConstVal, MIB2);
5447 MachineInstrBuilder MIB3 =
5448 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5451 GR.
add(ConstVal, MIB3);
5455 assert(NewReg != ResVReg);
5456 return BuildCOPY(ResVReg, NewReg,
I);
5466 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
5472 SPIRVTypeInst ResType =
5476 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
5481 if (
GlobalVar->isExternallyInitialized() &&
5482 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
5483 constexpr unsigned ReadWriteINTEL = 3u;
5486 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
5492bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
5493 SPIRVTypeInst ResType,
5494 MachineInstr &
I)
const {
5496 return selectExtInst(ResVReg, ResType,
I, CL::log10);
5504 MachineIRBuilder MIRBuilder(
I);
5509 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5512 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5514 .
add(
I.getOperand(1))
5519 ResType->
getOpcode() == SPIRV::OpTypeFloat);
5521 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
5529 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
5530 ? SPIRV::OpVectorTimesScalar
5541bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
5542 SPIRVTypeInst ResType,
5543 MachineInstr &
I)
const {
5559 MachineIRBuilder MIRBuilder(
I);
5562 ResType, MIRBuilder, SPIRV::StorageClass::Function);
5574 MachineBasicBlock &EntryBB =
I.getMF()->front();
5578 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
5581 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
5587 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5590 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
5593 .
add(
I.getOperand(
I.getNumExplicitDefs()))
5597 Register IntegralPartReg =
I.getOperand(1).getReg();
5598 if (IntegralPartReg.
isValid()) {
5600 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5611 assert(
false &&
"GLSL::Modf is deprecated.");
5622bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
5623 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5624 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5625 MachineIRBuilder MIRBuilder(
I);
5626 const SPIRVTypeInst Vec3Ty =
5629 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
5641 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5645 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
5646 Register LoadedRegister =
MRI->createVirtualRegister(&SPIRV::iIDRegClass);
5651 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5658 assert(
I.getOperand(2).isReg());
5659 const uint32_t ThreadId =
foldImm(
I.getOperand(2),
MRI);
5663 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5674bool SPIRVInstructionSelector::loadBuiltinInputID(
5675 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5676 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5677 MachineIRBuilder MIRBuilder(
I);
5679 ResType, MIRBuilder, SPIRV::StorageClass::Input);
5694 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5698 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5707SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
5708 MachineInstr &
I)
const {
5709 MachineIRBuilder MIRBuilder(
I);
5710 if (
Type->getOpcode() != SPIRV::OpTypeVector)
5714 if (VectorSize == 4)
5722bool SPIRVInstructionSelector::loadHandleBeforePosition(
5723 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
5724 MachineInstr &Pos)
const {
5727 Intrinsic::spv_resource_handlefrombinding);
5735 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
5736 MachineIRBuilder MIRBuilder(HandleDef);
5737 SPIRVTypeInst VarType = ResType;
5738 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
5740 if (IsStructuredBuffer) {
5746 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
5747 ArraySize, IndexReg, Name, MIRBuilder);
5751 uint32_t LoadOpcode =
5752 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
5762void SPIRVInstructionSelector::errorIfInstrOutsideShader(
5763 MachineInstr &
I)
const {
5765 std::string DiagMsg;
5766 raw_string_ostream OS(DiagMsg);
5767 I.print(OS,
true,
false,
false,
false);
5768 DiagMsg +=
" is only supported in shaders.\n";
5774InstructionSelector *
5778 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.
constexpr bool isScalar() const
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.
constexpr bool isPointer() const
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...