34#include "llvm/IR/IntrinsicsSPIRV.h"
40#define DEBUG_TYPE "spirv-isel"
47 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
52 std::optional<Register> Bias;
53 std::optional<Register>
Offset;
54 std::optional<Register> MinLod;
55 std::optional<Register> GradX;
56 std::optional<Register> GradY;
57 std::optional<Register> Lod;
58 std::optional<Register> Compare;
65 bool IsScalar =
false;
68llvm::SPIRV::SelectionControl::SelectionControl
69getSelectionOperandForImm(
int Imm) {
71 return SPIRV::SelectionControl::Flatten;
73 return SPIRV::SelectionControl::DontFlatten;
75 return SPIRV::SelectionControl::None;
79#define GET_GLOBALISEL_PREDICATE_BITSET
80#include "SPIRVGenGlobalISel.inc"
81#undef GET_GLOBALISEL_PREDICATE_BITSET
108#define GET_GLOBALISEL_PREDICATES_DECL
109#include "SPIRVGenGlobalISel.inc"
110#undef GET_GLOBALISEL_PREDICATES_DECL
112#define GET_GLOBALISEL_TEMPORARIES_DECL
113#include "SPIRVGenGlobalISel.inc"
114#undef GET_GLOBALISEL_TEMPORARIES_DECL
138 unsigned BitSetOpcode)
const;
142 unsigned BitSetOpcode)
const;
146 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
153 unsigned Opcode)
const;
156 unsigned Opcode)
const;
178 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
195 unsigned OpType)
const;
262 unsigned Opcode)
const;
266 unsigned Opcode)
const;
270 unsigned Opcode)
const;
274 unsigned Opcode)
const;
276 template <
bool Signed>
279 template <
bool Signed>
286 template <
typename PickOpcodeFn>
289 PickOpcodeFn &&PickOpcode)
const;
306 template <
typename PickOpcodeFn>
309 PickOpcodeFn &&PickOpcode)
const;
327 bool IsSigned)
const;
329 bool IsSigned,
unsigned Opcode)
const;
331 bool IsSigned)
const;
337 bool IsSigned)
const;
378 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
379 bool useMISrc =
true,
381 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
382 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
383 bool useMISrc =
true,
385 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
386 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
387 bool setMIFlags =
true,
bool useMISrc =
true,
389 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
390 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
391 bool useMISrc =
true,
394 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
395 MachineInstr &
I)
const;
397 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
398 MachineInstr &
I)
const;
400 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
401 MachineInstr &
I)
const;
403 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
404 MachineInstr &
I,
unsigned Opcode)
const;
406 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
407 bool WithGroupSync)
const;
409 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
412 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
417 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
418 MachineInstr &
I)
const;
420 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
421 MachineInstr &
I)
const;
423 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
424 MachineInstr &
I)
const;
425 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
426 MachineInstr &
I)
const;
427 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
428 SPIRVTypeInst ResType,
429 MachineInstr &
I)
const;
430 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
431 MachineInstr &
I)
const;
434 std::optional<Register> LodReg = std::nullopt)
const;
435 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
436 MachineInstr &
I)
const;
437 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
438 MachineInstr &
I)
const;
439 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
440 MachineInstr &
I)
const;
441 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
442 MachineInstr &
I)
const;
443 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
444 MachineInstr &
I)
const;
445 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
446 MachineInstr &
I)
const;
447 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
450 SPIRVTypeInst ResType,
451 MachineInstr &
I)
const;
452 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
453 MachineInstr &
I)
const;
454 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
455 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
458 MachineInstr &
I)
const;
459 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
460 MachineInstr &
I)
const;
461 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
462 MachineInstr &
I)
const;
463 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
464 MachineInstr &
I)
const;
465 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
466 MachineInstr &
I)
const;
467 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
468 MachineInstr &
I)
const;
469 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
470 MachineInstr &
I)
const;
471 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
472 MachineInstr &
I,
const unsigned DPdOpCode)
const;
474 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
475 SPIRVTypeInst ResType =
nullptr)
const;
476 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
477 SPIRVTypeInst ResType =
nullptr)
const;
479 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
480 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
481 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
483 MachineInstr &
I)
const;
484 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
486 bool wrapIntoSpecConstantOp(MachineInstr &
I,
489 Register getUcharPtrTypeReg(MachineInstr &
I,
490 SPIRV::StorageClass::StorageClass SC)
const;
491 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
493 uint32_t Opcode)
const;
494 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
495 SPIRVTypeInst SrcPtrTy)
const;
496 Register buildPointerToResource(SPIRVTypeInst ResType,
497 SPIRV::StorageClass::StorageClass SC,
498 uint32_t Set, uint32_t
Binding,
499 uint32_t ArraySize,
Register IndexReg,
501 MachineIRBuilder MIRBuilder)
const;
502 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
503 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
504 Register &ReadReg, MachineInstr &InsertionPoint)
const;
505 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
508 const ImageOperands *ImOps =
nullptr)
const;
509 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
511 Register CoordinateReg,
const ImageOperands &ImOps,
514 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
515 Register ResVReg, SPIRVTypeInst ResType,
516 MachineInstr &
I)
const;
517 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
518 Register ResVReg, SPIRVTypeInst ResType,
519 MachineInstr &
I)
const;
520 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
521 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
522 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
523 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
525 std::optional<SplitParts> splitEvenOddLanes(
Register PopCountReg,
526 unsigned ComponentCount,
528 SPIRVTypeInst I32Type)
const;
531 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
532 Register SrcReg,
unsigned int Opcode,
533 std::function<
bool(
Register, SPIRVTypeInst,
534 MachineInstr &,
Register,
unsigned)>
538bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
540 if (
TET->getTargetExtName() ==
"spirv.Image") {
543 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
544 return TET->getTypeParameter(0)->isIntegerTy();
548#define GET_GLOBALISEL_IMPL
549#include "SPIRVGenGlobalISel.inc"
550#undef GET_GLOBALISEL_IMPL
556 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
559#include
"SPIRVGenGlobalISel.inc"
562#include
"SPIRVGenGlobalISel.inc"
574 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
578void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
579 if (HasVRegsReset == &MF)
594 for (
const auto &
MBB : MF) {
595 for (
const auto &
MI :
MBB) {
598 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
602 LLT DstType = MRI.
getType(DstReg);
604 LLT SrcType = MRI.
getType(SrcReg);
605 if (DstType != SrcType)
610 if (DstRC != SrcRC && SrcRC)
622 while (!Stack.empty()) {
627 switch (
MI->getOpcode()) {
628 case TargetOpcode::G_INTRINSIC:
629 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
630 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
633 if (IntrID != Intrinsic::spv_const_composite &&
634 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
638 case TargetOpcode::G_BUILD_VECTOR:
639 case TargetOpcode::G_SPLAT_VECTOR:
641 i < OpDef->getNumOperands(); i++) {
646 Stack.push_back(OpNestedDef);
649 case TargetOpcode::G_CONSTANT:
650 case TargetOpcode::G_FCONSTANT:
651 case TargetOpcode::G_IMPLICIT_DEF:
652 case SPIRV::OpConstantTrue:
653 case SPIRV::OpConstantFalse:
654 case SPIRV::OpConstantI:
655 case SPIRV::OpConstantF:
656 case SPIRV::OpConstantComposite:
657 case SPIRV::OpConstantCompositeContinuedINTEL:
658 case SPIRV::OpConstantSampler:
659 case SPIRV::OpConstantNull:
661 case SPIRV::OpPoisonKHR:
662 case SPIRV::OpConstantFunctionPointerINTEL:
689 case Intrinsic::spv_all:
690 case Intrinsic::spv_alloca:
691 case Intrinsic::spv_any:
692 case Intrinsic::spv_bitcast:
693 case Intrinsic::spv_const_composite:
694 case Intrinsic::spv_cross:
695 case Intrinsic::spv_degrees:
696 case Intrinsic::spv_distance:
697 case Intrinsic::spv_extractelt:
698 case Intrinsic::spv_extractv:
699 case Intrinsic::spv_faceforward:
700 case Intrinsic::spv_fdot:
701 case Intrinsic::spv_firstbitlow:
702 case Intrinsic::spv_firstbitshigh:
703 case Intrinsic::spv_firstbituhigh:
704 case Intrinsic::spv_frac:
705 case Intrinsic::spv_gep:
706 case Intrinsic::spv_global_offset:
707 case Intrinsic::spv_global_size:
708 case Intrinsic::spv_group_id:
709 case Intrinsic::spv_insertelt:
710 case Intrinsic::spv_insertv:
711 case Intrinsic::spv_isinf:
712 case Intrinsic::spv_isnan:
713 case Intrinsic::spv_isfinite:
714 case Intrinsic::spv_isnormal:
715 case Intrinsic::spv_lerp:
716 case Intrinsic::spv_length:
717 case Intrinsic::spv_normalize:
718 case Intrinsic::spv_num_subgroups:
719 case Intrinsic::spv_num_workgroups:
720 case Intrinsic::spv_ptrcast:
721 case Intrinsic::spv_radians:
722 case Intrinsic::spv_reflect:
723 case Intrinsic::spv_refract:
724 case Intrinsic::spv_resource_getbasepointer:
725 case Intrinsic::spv_resource_getpointer:
726 case Intrinsic::spv_resource_handlefrombinding:
727 case Intrinsic::spv_resource_handlefromimplicitbinding:
728 case Intrinsic::spv_resource_nonuniformindex:
729 case Intrinsic::spv_resource_sample:
730 case Intrinsic::spv_rsqrt:
731 case Intrinsic::spv_saturate:
732 case Intrinsic::spv_sdot:
733 case Intrinsic::spv_sign:
734 case Intrinsic::spv_smoothstep:
735 case Intrinsic::spv_step:
736 case Intrinsic::spv_subgroup_id:
737 case Intrinsic::spv_subgroup_local_invocation_id:
738 case Intrinsic::spv_subgroup_max_size:
739 case Intrinsic::spv_subgroup_size:
740 case Intrinsic::spv_thread_id:
741 case Intrinsic::spv_thread_id_in_group:
742 case Intrinsic::spv_udot:
743 case Intrinsic::spv_undef:
744 case Intrinsic::spv_value_md:
745 case Intrinsic::spv_workgroup_size:
757 case SPIRV::OpTypeVoid:
758 case SPIRV::OpTypeBool:
759 case SPIRV::OpTypeInt:
760 case SPIRV::OpTypeFloat:
761 case SPIRV::OpTypeVector:
762 case SPIRV::OpTypeMatrix:
763 case SPIRV::OpTypeImage:
764 case SPIRV::OpTypeSampler:
765 case SPIRV::OpTypeSampledImage:
766 case SPIRV::OpTypeArray:
767 case SPIRV::OpTypeRuntimeArray:
768 case SPIRV::OpTypeStruct:
769 case SPIRV::OpTypeOpaque:
770 case SPIRV::OpTypePointer:
771 case SPIRV::OpTypeFunction:
772 case SPIRV::OpTypeEvent:
773 case SPIRV::OpTypeDeviceEvent:
774 case SPIRV::OpTypeReserveId:
775 case SPIRV::OpTypeQueue:
776 case SPIRV::OpTypePipe:
777 case SPIRV::OpTypeForwardPointer:
778 case SPIRV::OpTypePipeStorage:
779 case SPIRV::OpTypeNamedBarrier:
780 case SPIRV::OpTypeAccelerationStructureNV:
781 case SPIRV::OpTypeCooperativeMatrixNV:
782 case SPIRV::OpTypeCooperativeMatrixKHR:
792 if (
MI.getNumDefs() == 0)
795 for (
const auto &MO :
MI.all_defs()) {
797 if (
Reg.isPhysical()) {
802 if (
UseMI.getOpcode() != SPIRV::OpName) {
809 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
810 MI.isLifetimeMarker()) {
813 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
824 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
825 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
828 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
833 if (
MI.mayStore() ||
MI.isCall() ||
834 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
835 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
836 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
847 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
854void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
856 for (
const auto &MO :
MI.all_defs()) {
860 SmallVector<MachineInstr *, 4> UselessOpNames;
863 "There is still a use of the dead function.");
866 for (MachineInstr *OpNameMI : UselessOpNames) {
868 OpNameMI->eraseFromParent();
873void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
876 removeOpNamesForDeadMI(
MI);
877 MI.eraseFromParent();
880bool SPIRVInstructionSelector::select(MachineInstr &
I) {
881 resetVRegsType(*
I.getParent()->getParent());
883 assert(
I.getParent() &&
"Instruction should be in a basic block!");
884 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
889 removeDeadInstruction(
I);
896 if (Opcode == SPIRV::ASSIGN_TYPE) {
897 Register DstReg =
I.getOperand(0).getReg();
898 Register SrcReg =
I.getOperand(1).getReg();
901 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
902 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
903 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
904 Register SelectDstReg =
Def->getOperand(0).getReg();
905 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
907 assert(SuccessToSelectSelect);
909 Def->eraseFromParent();
916 bool Res = selectImpl(
I, *CoverageInfo);
918 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
919 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
923 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
935 }
else if (
I.getNumDefs() == 1) {
947 removeDeadInstruction(
I);
952 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
953 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
959 bool HasDefs =
I.getNumDefs() > 0;
962 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
963 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
964 if (spvSelect(ResVReg, ResType,
I)) {
966 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
977 case TargetOpcode::G_CONSTANT:
978 case TargetOpcode::G_FCONSTANT:
985 MachineInstr &
I)
const {
988 if (DstRC != SrcRC && SrcRC)
990 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
997bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
998 SPIRVTypeInst ResType,
999 MachineInstr &
I)
const {
1000 const unsigned Opcode =
I.getOpcode();
1002 return selectImpl(
I, *CoverageInfo);
1004 case TargetOpcode::G_CONSTANT:
1005 case TargetOpcode::G_FCONSTANT:
1006 return selectConst(ResVReg, ResType,
I);
1007 case TargetOpcode::G_GLOBAL_VALUE:
1008 return selectGlobalValue(ResVReg,
I);
1009 case TargetOpcode::G_IMPLICIT_DEF:
1010 return selectOpUndef(ResVReg, ResType,
I);
1011 case TargetOpcode::G_FREEZE:
1012 return selectFreeze(ResVReg, ResType,
I);
1014 case TargetOpcode::G_INTRINSIC:
1015 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1016 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1017 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1018 return selectIntrinsic(ResVReg, ResType,
I);
1019 case TargetOpcode::G_BITREVERSE:
1020 return selectBitreverse(ResVReg, ResType,
I);
1022 case TargetOpcode::G_BUILD_VECTOR:
1023 return selectBuildVector(ResVReg, ResType,
I);
1024 case TargetOpcode::G_SPLAT_VECTOR:
1025 return selectSplatVector(ResVReg, ResType,
I);
1026 case TargetOpcode::G_CONCAT_VECTORS:
1027 return selectConcatVectors(ResVReg, ResType,
I);
1029 case TargetOpcode::G_SHUFFLE_VECTOR: {
1030 MachineBasicBlock &BB = *
I.getParent();
1031 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1034 .
addUse(
I.getOperand(1).getReg())
1035 .
addUse(
I.getOperand(2).getReg());
1036 for (
auto V :
I.getOperand(3).getShuffleMask())
1041 case TargetOpcode::G_MEMMOVE:
1042 case TargetOpcode::G_MEMCPY:
1043 case TargetOpcode::G_MEMCPY_INLINE:
1044 case TargetOpcode::G_MEMSET:
1045 case TargetOpcode::G_MEMSET_INLINE:
1046 return selectMemOperation(ResVReg,
I);
1048 case TargetOpcode::G_ICMP:
1049 return selectICmp(ResVReg, ResType,
I);
1050 case TargetOpcode::G_FCMP:
1051 return selectFCmp(ResVReg, ResType,
I);
1053 case TargetOpcode::G_FRAME_INDEX:
1054 return selectFrameIndex(ResVReg, ResType,
I);
1056 case TargetOpcode::G_LOAD:
1057 return selectLoad(ResVReg, ResType,
I);
1058 case TargetOpcode::G_STORE:
1059 return selectStore(
I);
1061 case TargetOpcode::G_BR:
1062 return selectBranch(
I);
1063 case TargetOpcode::G_BRCOND:
1064 return selectBranchCond(
I);
1066 case TargetOpcode::G_PHI:
1067 return selectPhi(ResVReg,
I);
1069 case TargetOpcode::G_FPTOSI:
1070 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1071 case TargetOpcode::G_FPTOUI:
1072 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1074 case TargetOpcode::G_FPTOSI_SAT:
1075 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1076 case TargetOpcode::G_FPTOUI_SAT:
1077 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1079 case TargetOpcode::G_SITOFP:
1080 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1081 case TargetOpcode::G_UITOFP:
1082 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1084 case TargetOpcode::G_CTPOP:
1085 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1086 case TargetOpcode::G_SMIN:
1087 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1088 case TargetOpcode::G_UMIN:
1089 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1091 case TargetOpcode::G_SMAX:
1092 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1093 case TargetOpcode::G_UMAX:
1094 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1096 case TargetOpcode::G_SCMP:
1097 return selectSUCmp(ResVReg, ResType,
I,
true);
1098 case TargetOpcode::G_UCMP:
1099 return selectSUCmp(ResVReg, ResType,
I,
false);
1100 case TargetOpcode::G_LROUND:
1101 case TargetOpcode::G_LLROUND: {
1104 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1106 regForLround, *(
I.getParent()->getParent()));
1108 CL::round, GL::Round,
false);
1110 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1117 case TargetOpcode::G_STRICT_FMA:
1118 case TargetOpcode::G_FMA: {
1121 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1124 .
addUse(
I.getOperand(1).getReg())
1125 .
addUse(
I.getOperand(2).getReg())
1126 .
addUse(
I.getOperand(3).getReg())
1131 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1134 case TargetOpcode::G_STRICT_FLDEXP:
1135 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1137 case TargetOpcode::G_FPOW:
1138 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1139 case TargetOpcode::G_FPOWI:
1140 return selectFpowi(ResVReg, ResType,
I);
1142 case TargetOpcode::G_FEXP:
1143 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1144 case TargetOpcode::G_FEXP2:
1145 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1146 case TargetOpcode::G_FEXP10:
1147 return selectExp10(ResVReg, ResType,
I);
1149 case TargetOpcode::G_FMODF:
1150 return selectModf(ResVReg, ResType,
I);
1151 case TargetOpcode::G_FSINCOS:
1152 return selectSincos(ResVReg, ResType,
I);
1154 case TargetOpcode::G_FLOG:
1155 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1156 case TargetOpcode::G_FLOG2:
1157 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1158 case TargetOpcode::G_FLOG10:
1159 return selectLog10(ResVReg, ResType,
I);
1161 case TargetOpcode::G_FABS:
1162 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1163 case TargetOpcode::G_ABS:
1164 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1166 case TargetOpcode::G_FMINNUM:
1167 case TargetOpcode::G_FMINIMUM:
1168 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1169 case TargetOpcode::G_FMAXNUM:
1170 case TargetOpcode::G_FMAXIMUM:
1171 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1173 case TargetOpcode::G_FCOPYSIGN:
1174 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1176 case TargetOpcode::G_FCEIL:
1177 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1178 case TargetOpcode::G_FFLOOR:
1179 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1181 case TargetOpcode::G_FCOS:
1182 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1183 case TargetOpcode::G_FSIN:
1184 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1185 case TargetOpcode::G_FTAN:
1186 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1187 case TargetOpcode::G_FACOS:
1188 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1189 case TargetOpcode::G_FASIN:
1190 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1191 case TargetOpcode::G_FATAN:
1192 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1193 case TargetOpcode::G_FATAN2:
1194 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1195 case TargetOpcode::G_FCOSH:
1196 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1197 case TargetOpcode::G_FSINH:
1198 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1199 case TargetOpcode::G_FTANH:
1200 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1202 case TargetOpcode::G_STRICT_FSQRT:
1203 case TargetOpcode::G_FSQRT:
1204 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1206 case TargetOpcode::G_CTTZ:
1207 case TargetOpcode::G_CTTZ_ZERO_POISON:
1208 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1209 case TargetOpcode::G_CTLZ:
1210 case TargetOpcode::G_CTLZ_ZERO_POISON:
1211 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1213 case TargetOpcode::G_INTRINSIC_ROUND:
1214 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1215 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1216 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1217 case TargetOpcode::G_INTRINSIC_TRUNC:
1218 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1219 case TargetOpcode::G_FRINT:
1220 case TargetOpcode::G_FNEARBYINT:
1221 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1223 case TargetOpcode::G_SMULH:
1224 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1225 case TargetOpcode::G_UMULH:
1226 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1228 case TargetOpcode::G_SADDSAT:
1229 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1230 case TargetOpcode::G_UADDSAT:
1231 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1232 case TargetOpcode::G_SSUBSAT:
1233 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1234 case TargetOpcode::G_USUBSAT:
1235 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1237 case TargetOpcode::G_FFREXP:
1238 return selectFrexp(ResVReg, ResType,
I);
1240 case TargetOpcode::G_UADDO:
1241 return selectOverflowArith(ResVReg, ResType,
I,
1242 ResType->
getOpcode() == SPIRV::OpTypeVector
1243 ? SPIRV::OpIAddCarryV
1244 : SPIRV::OpIAddCarryS);
1245 case TargetOpcode::G_USUBO:
1246 return selectOverflowArith(ResVReg, ResType,
I,
1247 ResType->
getOpcode() == SPIRV::OpTypeVector
1248 ? SPIRV::OpISubBorrowV
1249 : SPIRV::OpISubBorrowS);
1250 case TargetOpcode::G_UMULO:
1251 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1252 case TargetOpcode::G_SMULO:
1253 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1255 case TargetOpcode::G_SEXT:
1256 return selectExt(ResVReg, ResType,
I,
true);
1257 case TargetOpcode::G_ANYEXT:
1258 case TargetOpcode::G_ZEXT:
1259 return selectExt(ResVReg, ResType,
I,
false);
1260 case TargetOpcode::G_TRUNC:
1261 return selectTrunc(ResVReg, ResType,
I);
1262 case TargetOpcode::G_FPTRUNC:
1263 case TargetOpcode::G_FPEXT:
1264 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1266 case TargetOpcode::G_PTRTOINT:
1267 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1268 case TargetOpcode::G_INTTOPTR:
1269 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1270 case TargetOpcode::G_BITCAST:
1271 return selectBitcast(ResVReg, ResType,
I);
1272 case TargetOpcode::G_ADDRSPACE_CAST:
1273 return selectAddrSpaceCast(ResVReg, ResType,
I);
1274 case TargetOpcode::G_PTRMASK:
1275 return selectPtrMask(ResVReg, ResType,
I);
1276 case TargetOpcode::G_PTR_ADD: {
1278 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1282 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1283 (*II).getOpcode() == TargetOpcode::COPY ||
1284 (*II).getOpcode() == SPIRV::OpVariable) &&
1285 getImm(
I.getOperand(2), MRI));
1287 bool IsGVInit =
false;
1291 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1292 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1293 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1294 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1304 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1316 return diagnoseUnsupported(
1317 I,
"incompatible result and operand types in a bitcast");
1319 MachineInstrBuilder MIB =
1320 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1327 : SPIRV::OpInBoundsPtrAccessChain))
1331 .
addUse(
I.getOperand(2).getReg())
1334 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1338 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1340 .
addUse(
I.getOperand(2).getReg())
1349 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1352 .
addImm(
static_cast<uint32_t
>(
1353 SPIRV::Opcode::InBoundsPtrAccessChain))
1356 .
addUse(
I.getOperand(2).getReg());
1361 case TargetOpcode::G_ATOMICRMW_OR:
1362 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1363 case TargetOpcode::G_ATOMICRMW_ADD:
1364 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1365 case TargetOpcode::G_ATOMICRMW_AND:
1366 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1367 case TargetOpcode::G_ATOMICRMW_MAX:
1368 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1369 case TargetOpcode::G_ATOMICRMW_MIN:
1370 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1371 case TargetOpcode::G_ATOMICRMW_SUB:
1372 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1373 case TargetOpcode::G_ATOMICRMW_XOR:
1374 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1375 case TargetOpcode::G_ATOMICRMW_UMAX:
1376 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1377 case TargetOpcode::G_ATOMICRMW_UMIN:
1378 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1379 case TargetOpcode::G_ATOMICRMW_XCHG:
1380 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1382 case TargetOpcode::G_ATOMICRMW_FADD:
1383 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1384 case TargetOpcode::G_ATOMICRMW_FSUB:
1386 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1387 ResType->
getOpcode() == SPIRV::OpTypeVector
1389 : SPIRV::OpFNegate);
1390 case TargetOpcode::G_ATOMICRMW_FMIN:
1391 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1392 case TargetOpcode::G_ATOMICRMW_FMAX:
1393 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1395 case TargetOpcode::G_FENCE:
1396 return selectFence(
I);
1398 case TargetOpcode::G_STACKSAVE:
1399 return selectStackSave(ResVReg, ResType,
I);
1400 case TargetOpcode::G_STACKRESTORE:
1401 return selectStackRestore(
I);
1403 case TargetOpcode::G_UNMERGE_VALUES:
1406 case TargetOpcode::G_TRAP:
1407 case TargetOpcode::G_UBSANTRAP:
1408 return selectTrap(
I);
1413 case TargetOpcode::DBG_LABEL:
1415 case TargetOpcode::G_DEBUGTRAP:
1416 return selectDebugTrap(ResVReg, ResType,
I);
1423bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1424 SPIRVTypeInst ResType,
1425 MachineInstr &
I)
const {
1426 unsigned Opcode = SPIRV::OpNop;
1433bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1434 SPIRVTypeInst ResType,
1436 GL::GLSLExtInst GLInst,
1437 bool setMIFlags,
bool useMISrc,
1440 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1441 return diagnoseUnsupported(
1443 "this instruction is only supported with the GLSL extended instruction "
1445 return selectExtInst(ResVReg, ResType,
I,
1446 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1447 setMIFlags, useMISrc, SrcRegs);
1450bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1451 SPIRVTypeInst ResType,
1453 CL::OpenCLExtInst CLInst,
1454 bool setMIFlags,
bool useMISrc,
1456 return selectExtInst(ResVReg, ResType,
I,
1457 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1458 setMIFlags, useMISrc, SrcRegs);
1461bool SPIRVInstructionSelector::selectExtInst(
1462 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1463 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1465 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1466 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1467 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1471bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1472 SPIRVTypeInst ResType,
1475 bool setMIFlags,
bool useMISrc,
1478 for (
const auto &[InstructionSet, Opcode] : Insts) {
1482 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1485 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1490 const unsigned NumOps =
I.getNumOperands();
1493 I.getOperand(Index).getType() ==
1494 MachineOperand::MachineOperandType::MO_IntrinsicID)
1497 MIB.
add(
I.getOperand(Index));
1509bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1510 SPIRVTypeInst ResType,
1511 MachineInstr &
I)
const {
1512 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1513 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1514 for (
const auto &Ex : ExtInsts) {
1515 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1516 uint32_t Opcode = Ex.second;
1520 MachineIRBuilder MIRBuilder(
I);
1523 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1528 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1531 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1535 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1538 .
addImm(
static_cast<uint32_t
>(Ex.first))
1540 .
add(
I.getOperand(2))
1544 Register ExpResReg =
I.getOperand(1).getReg();
1546 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1556bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1557 SPIRVTypeInst ResType,
1558 MachineInstr &
I)
const {
1559 Register CosResVReg =
I.getOperand(1).getReg();
1560 unsigned SrcIdx =
I.getNumExplicitDefs();
1565 MachineIRBuilder MIRBuilder(
I);
1567 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1572 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1575 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1577 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1580 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1582 .
add(
I.getOperand(SrcIdx))
1585 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1593 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1596 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1598 .
add(
I.getOperand(SrcIdx))
1600 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1603 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1605 .
add(
I.getOperand(SrcIdx))
1612bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1613 SPIRVTypeInst ResType,
1616 unsigned Opcode)
const {
1617 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1627std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1628 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1629 SPIRVTypeInst I32Type)
const {
1632 if (ComponentCount == 1) {
1635 Parts.IsScalar =
true;
1636 Parts.Type = I32Type;
1644 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1645 SPIRV::OpVectorExtractDynamic))
1646 return std::nullopt;
1648 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1649 SPIRV::OpVectorExtractDynamic))
1650 return std::nullopt;
1654 MachineIRBuilder MIRBuilder(
I);
1655 Parts.IsScalar =
false;
1662 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1663 TII.get(SPIRV::OpVectorShuffle))
1668 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1673 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1674 TII.get(SPIRV::OpVectorShuffle))
1679 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1687bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1688 SPIRVTypeInst ResType,
1691 unsigned Opcode)
const {
1692 Register OpReg =
I.getOperand(1).getReg();
1695 MachineIRBuilder MIRBuilder(
I);
1697 SPIRVTypeInst I32VectorType =
1700 bool IsVector = NumElems > 1;
1701 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1704 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1708 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1711 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1714bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1715 SPIRVTypeInst ResType,
1718 unsigned Opcode)
const {
1719 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1722bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1723 SPIRVTypeInst ResType,
1726 unsigned Opcode)
const {
1728 if (ComponentCount > 2)
1729 return handle64BitOverflow(
1730 ResVReg, ResType,
I, SrcReg, Opcode,
1732 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1734 MachineIRBuilder MIRBuilder(
I);
1739 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1743 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1748 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1752 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1755 SplitParts &Parts = *MaybeParts;
1758 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1760 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1765 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1766 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1769bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1770 SPIRVTypeInst ResType,
1772 unsigned Opcode)
const {
1777 if (!STI.getTargetTriple().isVulkanOS())
1778 return selectUnOp(ResVReg, ResType,
I, Opcode);
1780 Register OpReg =
I.getOperand(1).getReg();
1783 : SPIRV::OpUConvert;
1787 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1789 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1791 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1793 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1797bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1798 SPIRVTypeInst ResType,
1800 unsigned Opcode)
const {
1802 Register SrcReg =
I.getOperand(1).getReg();
1807 unsigned DefOpCode = DefIt->getOpcode();
1808 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1811 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1812 DefOpCode = VRD->getOpcode();
1814 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1815 DefOpCode == TargetOpcode::G_CONSTANT ||
1816 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1822 uint32_t SpecOpcode = 0;
1824 case SPIRV::OpConvertPtrToU:
1825 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1827 case SPIRV::OpConvertUToPtr:
1828 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1833 TII.get(SPIRV::OpSpecConstantOp))
1843 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1847bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1848 SPIRVTypeInst ResType,
1849 MachineInstr &
I)
const {
1850 Register OpReg =
I.getOperand(1).getReg();
1851 SPIRVTypeInst OpType =
1854 return diagnoseUnsupported(
1855 I,
"incompatible result and operand types in a bitcast");
1856 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1867 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1868 if (
MemOp->isNonTemporal())
1869 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1871 if (!ST->isShader() &&
MemOp->getAlign().value())
1872 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1876 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1877 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1881 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1883 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1887 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1891 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1893 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1905 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1907 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1909 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1913bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1914 SPIRVTypeInst ResType,
1915 MachineInstr &
I)
const {
1917 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1922 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1923 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1925 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1927 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1931 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1935 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1936 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1937 I.getDebugLoc(),
I);
1941 MachineIRBuilder MIRBuilder(
I);
1943 if (
I.getNumMemOperands()) {
1944 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1945 if (MemOp->isAtomic())
1946 return selectAtomicLoad(ResVReg, ResType,
I);
1949 auto MIB = MIRBuilder.buildInstr(SPIRV::OpLoad)
1953 if (!
I.getNumMemOperands()) {
1954 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1956 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1965bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
1966 SPIRVTypeInst ResType,
1967 MachineInstr &
I)
const {
1968 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
1971 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1974 return diagnoseUnsupported(
1975 I,
"Lowering to SPIR-V of atomic load is only "
1976 "allowed for integer, floating point or pointer types");
1978 assert(
I.getNumMemOperands());
1979 const MachineMemOperand &MemOp = **
I.memoperands_begin();
1980 assert(MemOp.isAtomic());
1984 Register ScopeReg = buildI32Constant(Scope,
I);
1990 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
1991 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
1994 MachineIRBuilder MIRBuilder(
I);
1998 return diagnoseUnsupported(
1999 I,
"Lowering to SPIR-V of atomic load is only "
2000 "allowed for pointer types for physical addressing model");
2007 SPIRVTypeInst PtrAsIntSpirvType =
2018 PtrAsIntSpirvType, MIRBuilder,
2021 MIRBuilder.getMF());
2023 MIRBuilder.buildInstr(SPIRV::OpBitcast)
2024 .addDef(PtrCastedToMatchValReg)
2027 .constrainAllUses(
TII,
TRI, RBI);
2029 MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
2032 .addUse(PtrCastedToMatchValReg)
2035 .constrainAllUses(
TII,
TRI, RBI);
2036 MIRBuilder.buildInstr(SPIRV::OpConvertUToPtr)
2040 .constrainAllUses(
TII,
TRI, RBI);
2043 auto AtomicLoad = MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
2049 AtomicLoad.constrainAllUses(
TII,
TRI, RBI);
2054bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2056 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2057 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2062 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2063 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2065 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2070 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2074 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2075 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2076 SPIRVTypeInst SampledType =
2078 SPIRVTypeInst StoreValCompType =
2080 if (StoreValCompType && StoreValCompType != SampledType) {
2083 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2086 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2091 StoreVal = PackedReg;
2094 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2095 TII.get(SPIRV::OpImageWrite))
2101 if (sampledTypeIsSignedInteger(LLVMHandleType))
2104 BMI.constrainAllUses(
TII,
TRI, RBI);
2109 if (
I.getNumMemOperands()) {
2110 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2111 if (MemOp->isAtomic())
2112 return selectAtomicStore(
I);
2119 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2120 PtrSC == SPIRV::StorageClass::Input ||
2121 PtrSC == SPIRV::StorageClass::PushConstant)
2122 return diagnoseUnsupported(
2123 I,
"store into a read-only SPIR-V storage class is not allowed");
2125 MachineIRBuilder MIRBuilder(
I);
2126 auto MIB = MIRBuilder.buildInstr(SPIRV::OpStore).
addUse(Ptr).
addUse(StoreVal);
2127 if (!
I.getNumMemOperands()) {
2128 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2130 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2139bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2140 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2143 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2144 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2149 assert(
I.getNumMemOperands());
2150 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2151 assert(MemOp.isAtomic());
2155 Register ScopeReg = buildI32Constant(Scope,
I);
2161 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2162 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2164 MachineIRBuilder MIRBuilder(
I);
2168 return diagnoseUnsupported(
2169 I,
"Lowering to SPIR-V of atomic store is only "
2170 "allowed for pointer types for physical addressing model");
2176 SPIRVTypeInst PtrAsIntSpirvType =
2183 MIRBuilder.buildInstr(SPIRV::OpConvertPtrToU)
2187 .constrainAllUses(
TII,
TRI, RBI);
2193 PtrAsIntSpirvType, MIRBuilder,
2196 MIRBuilder.getMF());
2198 MIRBuilder.buildInstr(SPIRV::OpBitcast)
2199 .addDef(PtrCastedToMatchValReg)
2202 .constrainAllUses(
TII,
TRI, RBI);
2204 StoreVal = PtrToUVal;
2205 Ptr = PtrCastedToMatchValReg;
2206 PointeeType = PtrAsIntSpirvType;
2210 return diagnoseUnsupported(
I,
2211 "Lowering to SPIR-V of atomic store is only "
2212 "allowed for integer or floating point types");
2214 auto AtomicStore = MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
2219 AtomicStore.constrainAllUses(
TII,
TRI, RBI);
2224bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2225 SPIRVTypeInst ResType,
2226 MachineInstr &
I)
const {
2227 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2235 const Register PtrsReg =
I.getOperand(2).getReg();
2236 const uint32_t Alignment =
I.getOperand(3).getImm();
2237 const Register MaskReg =
I.getOperand(4).getReg();
2238 const Register PassthruReg =
I.getOperand(5).getReg();
2239 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2243 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2254bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2255 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2262 const Register ValuesReg =
I.getOperand(1).getReg();
2263 const Register PtrsReg =
I.getOperand(2).getReg();
2264 const uint32_t Alignment =
I.getOperand(3).getImm();
2265 const Register MaskReg =
I.getOperand(4).getReg();
2266 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2270 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2279bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2280 const Twine &
Msg)
const {
2281 const Function &
F =
I.getMF()->getFunction();
2282 F.getContext().diagnose(
2283 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2287bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2288 SPIRVTypeInst ResType,
2289 MachineInstr &
I)
const {
2290 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2291 return diagnoseUnsupported(
2292 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2293 "SPIR-V extension: SPV_INTEL_variable_length_array");
2295 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2302bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2303 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2304 return diagnoseUnsupported(
2306 "llvm.stackrestore intrinsic: this instruction requires the following "
2307 "SPIR-V extension: SPV_INTEL_variable_length_array");
2308 if (!
I.getOperand(0).isReg())
2311 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2312 .
addUse(
I.getOperand(0).getReg())
2318SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2319 MachineIRBuilder MIRBuilder(
I);
2320 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2327 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2331 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2332 Type *ArrTy = ArrayType::get(ValTy, Num);
2334 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2337 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2344 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
2347 .
addImm(SPIRV::StorageClass::UniformConstant)
2358bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2361 Register DstReg =
I.getOperand(0).getReg();
2365 return diagnoseUnsupported(
2366 I,
"OpCopyMemory requires operands to have the same type");
2367 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2371 return diagnoseUnsupported(
2372 I,
"Unable to determine pointee type size for OpCopyMemory");
2373 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2374 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2375 return diagnoseUnsupported(
2376 I,
"OpCopyMemory requires the size to match the pointee type size");
2377 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2380 if (
I.getNumMemOperands()) {
2381 MachineIRBuilder MIRBuilder(
I);
2388bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2391 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2392 .
addUse(
I.getOperand(0).getReg())
2394 .
addUse(
I.getOperand(2).getReg());
2395 if (
I.getNumMemOperands()) {
2396 MachineIRBuilder MIRBuilder(
I);
2403bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2404 MachineInstr &
I)
const {
2406 Register SizeReg =
I.getOperand(2).getReg();
2408 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2412 Register SrcReg =
I.getOperand(1).getReg();
2413 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2414 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2415 Register VarReg = getOrCreateMemSetGlobal(
I);
2418 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2420 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2422 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2426 if (!selectCopyMemory(
I, SrcReg))
2429 if (!selectCopyMemorySized(
I, SrcReg))
2432 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2433 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2438bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2439 SPIRVTypeInst ResType,
2442 unsigned NegateOpcode)
const {
2444 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2447 Register ScopeReg = buildI32Constant(Scope,
I);
2449 Register Ptr =
I.getOperand(1).getReg();
2450 uint32_t ScSem =
static_cast<uint32_t
>(
2454 Register MemSemReg = buildI32Constant(MemSem,
I);
2456 Register ValueReg =
I.getOperand(2).getReg();
2457 if (NegateOpcode != 0) {
2460 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2466 if (NewOpcode != SPIRV::OpAtomicExchange)
2467 return diagnoseUnsupported(
2468 I,
"Lowering to SPIR-V of this atomic operation is not "
2469 "allowed for pointer types");
2471 return diagnoseUnsupported(
2472 I,
"Lowering to SPIR-V of atomic exchange is only "
2473 "allowed for pointer types for physical addressing model");
2480 MachineIRBuilder MIRBuilder(
I);
2482 SPIRVTypeInst PtrAsIntSpirvType =
2489 MIRBuilder.getMF());
2490 MIRBuilder.buildInstr(SPIRV::OpConvertPtrToU)
2491 .addDef(ValueAsIntReg)
2494 .constrainAllUses(
TII,
TRI, RBI);
2502 MIRBuilder.getMF());
2503 MIRBuilder.buildInstr(SPIRV::OpBitcast)
2504 .addDef(PtrCastedToMatchValReg)
2507 .constrainAllUses(
TII,
TRI, RBI);
2513 MIRBuilder.getMF());
2514 MIRBuilder.buildInstr(SPIRV::OpAtomicExchange)
2515 .addDef(ExchangeResReg)
2517 .addUse(PtrCastedToMatchValReg)
2520 .addUse(ValueAsIntReg)
2521 .constrainAllUses(
TII,
TRI, RBI);
2522 MIRBuilder.buildInstr(SPIRV::OpConvertUToPtr)
2525 .addUse(ExchangeResReg)
2526 .constrainAllUses(
TII,
TRI, RBI);
2530 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2541bool SPIRVInstructionSelector::selectInterlockedOp(
Register ResVReg,
2542 SPIRVTypeInst ResType,
2544 unsigned Opcode)
const {
2545 Register Ptr =
I.getOperand(2).getReg();
2549 assert((SC == SPIRV::StorageClass::Workgroup ||
2550 SC == SPIRV::StorageClass::StorageBuffer) &&
2551 "InterlockedAdd requires Workgroup or StorageBuffer storage class");
2552 uint32_t
Scope =
static_cast<uint32_t
>(SC == SPIRV::StorageClass::Workgroup
2553 ? SPIRV::Scope::Workgroup
2554 : SPIRV::Scope::Device);
2555 Register ScopeReg = buildI32Constant(Scope,
I);
2558 Register MemSemReg = buildI32Constant(MemSem,
I);
2560 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
2571bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2572 unsigned ArgI =
I.getNumOperands() - 1;
2574 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2575 SPIRVTypeInst SrcType =
2577 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2579 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2583 unsigned CurrentIndex = 0;
2584 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2585 Register ResVReg =
I.getOperand(i).getReg();
2588 LLT ResLLT = MRI->
getType(ResVReg);
2594 ResType = ScalarType;
2600 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2603 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2609 for (
unsigned j = 0;
j < NumElements; ++
j) {
2610 MIB.
addImm(CurrentIndex + j);
2612 CurrentIndex += NumElements;
2616 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2628bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2631 Register MemSemReg = buildI32Constant(MemSem,
I);
2633 uint32_t
Scope =
static_cast<uint32_t
>(
2635 Register ScopeReg = buildI32Constant(Scope,
I);
2637 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2644bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2645 SPIRVTypeInst ResType,
2647 unsigned Opcode)
const {
2648 Type *ResTy =
nullptr;
2651 return diagnoseUnsupported(
2653 "Not enough info to select the arithmetic with overflow instruction");
2655 return diagnoseUnsupported(
I,
2656 "Expect struct type result for the arithmetic "
2657 "with overflow instruction");
2663 MachineIRBuilder MIRBuilder(
I);
2665 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2666 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2672 Register ZeroReg = buildZerosVal(ResType,
I);
2677 if (ResName.
size() > 0)
2682 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2685 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2686 MIB.
addUse(
I.getOperand(i).getReg());
2691 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2692 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2694 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2695 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2702 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2703 .
addDef(
I.getOperand(1).getReg())
2711bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2712 SPIRVTypeInst ResType,
2713 MachineInstr &
I)
const {
2715 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2716 Register Ptr =
I.getOperand(2).getReg();
2717 Register ScopeReg =
I.getOperand(5).getReg();
2718 Register MemSemEqReg =
I.getOperand(6).getReg();
2719 Register MemSemNeqReg =
I.getOperand(7).getReg();
2721 Register Val =
I.getOperand(4).getReg();
2725 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2744 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2751 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2763 case SPIRV::StorageClass::DeviceOnlyINTEL:
2764 case SPIRV::StorageClass::HostOnlyINTEL:
2773 bool IsGRef =
false;
2774 bool IsAllowedRefs =
2776 unsigned Opcode = It.getOpcode();
2777 if (Opcode == SPIRV::OpConstantComposite ||
2778 Opcode == SPIRV::OpSpecConstantComposite ||
2779 Opcode == SPIRV::OpVariable ||
2780 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2781 return IsGRef = true;
2782 return Opcode == SPIRV::OpName;
2784 return IsAllowedRefs && IsGRef;
2787Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2788 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2790 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2794SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2796 uint32_t Opcode)
const {
2797 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2798 TII.get(SPIRV::OpSpecConstantOp))
2806SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2807 SPIRVTypeInst SrcPtrTy)
const {
2808 SPIRVTypeInst GenericPtrTy =
2812 SPIRV::StorageClass::Generic),
2814 MachineFunction *MF =
I.getParent()->getParent();
2816 MachineInstrBuilder MIB = buildSpecConstantOp(
2818 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2828bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2829 SPIRVTypeInst ResType,
2830 MachineInstr &
I)
const {
2834 Register SrcPtr =
I.getOperand(1).getReg();
2838 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2839 ResType->
getOpcode() != SPIRV::OpTypePointer)
2840 return BuildCOPY(ResVReg, SrcPtr,
I);
2850 unsigned SpecOpcode =
2852 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2855 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2862 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2864 .constrainAllUses(
TII,
TRI, RBI);
2866 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2868 buildSpecConstantOp(
2870 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2871 .constrainAllUses(
TII,
TRI, RBI);
2878 return BuildCOPY(ResVReg, SrcPtr,
I);
2880 if ((SrcSC == SPIRV::StorageClass::Function &&
2881 DstSC == SPIRV::StorageClass::Private) ||
2882 (DstSC == SPIRV::StorageClass::Function &&
2883 SrcSC == SPIRV::StorageClass::Private))
2884 return BuildCOPY(ResVReg, SrcPtr,
I);
2888 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2891 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2894 SPIRVTypeInst GenericPtrTy =
2913 return selectUnOp(ResVReg, ResType,
I,
2914 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2916 return selectUnOp(ResVReg, ResType,
I,
2917 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2919 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2921 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2931bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
2932 SPIRVTypeInst ResType,
2933 MachineInstr &
I)
const {
2935 return diagnoseUnsupported(
2936 I,
"G_PTRMASK is not supported with logical SPIR-V");
2941 Register PtrReg =
I.getOperand(1).getReg();
2942 Register MaskReg =
I.getOperand(2).getReg();
2961 ? SPIRV::OpBitwiseAndV
2962 : SPIRV::OpBitwiseAndS;
2985 return SPIRV::OpFOrdEqual;
2987 return SPIRV::OpFOrdGreaterThanEqual;
2989 return SPIRV::OpFOrdGreaterThan;
2991 return SPIRV::OpFOrdLessThanEqual;
2993 return SPIRV::OpFOrdLessThan;
2995 return SPIRV::OpFOrdNotEqual;
2997 return SPIRV::OpOrdered;
2999 return SPIRV::OpFUnordEqual;
3001 return SPIRV::OpFUnordGreaterThanEqual;
3003 return SPIRV::OpFUnordGreaterThan;
3005 return SPIRV::OpFUnordLessThanEqual;
3007 return SPIRV::OpFUnordLessThan;
3009 return SPIRV::OpFUnordNotEqual;
3011 return SPIRV::OpUnordered;
3021 return SPIRV::OpIEqual;
3023 return SPIRV::OpINotEqual;
3025 return SPIRV::OpSGreaterThanEqual;
3027 return SPIRV::OpSGreaterThan;
3029 return SPIRV::OpSLessThanEqual;
3031 return SPIRV::OpSLessThan;
3033 return SPIRV::OpUGreaterThanEqual;
3035 return SPIRV::OpUGreaterThan;
3037 return SPIRV::OpULessThanEqual;
3039 return SPIRV::OpULessThan;
3048 return SPIRV::OpPtrEqual;
3050 return SPIRV::OpPtrNotEqual;
3061 return SPIRV::OpLogicalEqual;
3063 return SPIRV::OpLogicalNotEqual;
3101bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3102 SPIRVTypeInst ResType,
3104 unsigned OpAnyOrAll)
const {
3105 assert(
I.getNumOperands() == 3);
3106 assert(
I.getOperand(2).isReg());
3108 Register InputRegister =
I.getOperand(2).getReg();
3111 assert(InputType &&
"VReg has no type assigned");
3114 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
3115 if (IsBoolTy && !IsVectorTy) {
3116 assert(ResVReg ==
I.getOperand(0).getReg());
3117 return BuildCOPY(ResVReg, InputRegister,
I);
3121 unsigned SpirvNotEqualId =
3122 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3124 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3129 IsBoolTy ? InputRegister
3137 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3139 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3156bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3157 SPIRVTypeInst ResType,
3158 MachineInstr &
I)
const {
3159 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3162bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3163 SPIRVTypeInst ResType,
3164 MachineInstr &
I)
const {
3165 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3169bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3170 SPIRVTypeInst ResType,
3171 MachineInstr &
I)
const {
3172 assert(
I.getNumOperands() == 4);
3173 assert(
I.getOperand(2).isReg());
3174 assert(
I.getOperand(3).isReg());
3176 [[maybe_unused]] SPIRVTypeInst VecType =
3181 "dot product requires a vector of at least 2 components");
3183 [[maybe_unused]] SPIRVTypeInst EltType =
3192 .
addUse(
I.getOperand(2).getReg())
3193 .
addUse(
I.getOperand(3).getReg())
3198bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3199 SPIRVTypeInst ResType,
3202 assert(
I.getNumOperands() == 4);
3203 assert(
I.getOperand(2).isReg());
3204 assert(
I.getOperand(3).isReg());
3207 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3211 .
addUse(
I.getOperand(2).getReg())
3212 .
addUse(
I.getOperand(3).getReg())
3219bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3220 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3221 assert(
I.getNumOperands() == 4);
3222 assert(
I.getOperand(2).isReg());
3223 assert(
I.getOperand(3).isReg());
3227 Register Vec0 =
I.getOperand(2).getReg();
3228 Register Vec1 =
I.getOperand(3).getReg();
3232 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3241 "dot product requires a vector of at least 2 components");
3244 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3254 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3265 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3277bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3278 SPIRVTypeInst ResType,
3279 MachineInstr &
I)
const {
3281 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3284 .
addUse(
I.getOperand(2).getReg())
3289bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3290 SPIRVTypeInst ResType,
3291 MachineInstr &
I)
const {
3293 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3296 .
addUse(
I.getOperand(2).getReg())
3301bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3302 SPIRVTypeInst ResType,
3303 MachineInstr &
I)
const {
3305 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3308 .
addUse(
I.getOperand(2).getReg())
3313bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3314 SPIRVTypeInst ResType,
3315 MachineInstr &
I)
const {
3317 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3320 .
addUse(
I.getOperand(2).getReg())
3325template <
bool Signed>
3326bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3327 SPIRVTypeInst ResType,
3328 MachineInstr &
I)
const {
3329 assert(
I.getNumOperands() == 5);
3330 assert(
I.getOperand(2).isReg());
3331 assert(
I.getOperand(3).isReg());
3332 assert(
I.getOperand(4).isReg());
3335 Register Acc =
I.getOperand(2).getReg();
3339 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3341 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3346 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3349 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3361template <
bool Signed>
3362bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3363 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3364 assert(
I.getNumOperands() == 5);
3365 assert(
I.getOperand(2).isReg());
3366 assert(
I.getOperand(3).isReg());
3367 assert(
I.getOperand(4).isReg());
3370 Register Acc =
I.getOperand(2).getReg();
3376 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3380 for (
unsigned i = 0; i < 4; i++) {
3403 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3423 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3438bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3439 SPIRVTypeInst ResType,
3440 MachineInstr &
I)
const {
3441 assert(
I.getNumOperands() == 3);
3442 assert(
I.getOperand(2).isReg());
3444 Register VZero = buildZerosValF(ResType,
I);
3445 Register VOne = buildOnesValF(ResType,
I);
3447 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3450 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3452 .
addUse(
I.getOperand(2).getReg())
3459bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3460 SPIRVTypeInst ResType,
3461 MachineInstr &
I)
const {
3462 assert(
I.getNumOperands() == 3);
3463 assert(
I.getOperand(2).isReg());
3465 Register InputRegister =
I.getOperand(2).getReg();
3467 auto &
DL =
I.getDebugLoc();
3470 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3477 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3479 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3487 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3492 if (NeedsConversion) {
3493 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3504bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3505 SPIRVTypeInst ResType,
3507 unsigned Opcode)
const {
3511 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3517 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3518 BMI.addUse(
I.getOperand(J).getReg());
3525bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3528 bool WithGroupSync)
const {
3530 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3532 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3534 assert(((Scope != SPIRV::Scope::Workgroup) ||
3535 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3536 "Workgroup Scope must set WorkGroupMemory semantic "
3537 "in Barrier instruction");
3539 assert(((Scope != SPIRV::Scope::Device) ||
3540 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3541 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3542 "Device Scope must set UniformMemory and ImageMemory semantic "
3543 "in Barrier instruction");
3549 if (WithGroupSync) {
3550 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3554 Register ScopeReg = buildI32Constant(Scope,
I);
3555 Register MemSemReg = buildI32Constant(MemSem,
I);
3557 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3561bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3562 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3567 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3568 SPIRV::OpGroupNonUniformBallot))
3573 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3578 .
addImm(SPIRV::GroupOperation::Reduce)
3585bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3586 SPIRVTypeInst ResType,
3587 MachineInstr &
I)
const {
3592 Register InputReg =
I.getOperand(2).getReg();
3597 bool IsVector = NumElems > 1;
3610 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3611 SPIRV::OpGroupNonUniformAllEqual);
3616 ElementResults.
reserve(NumElems);
3618 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3631 ElemInput = Extracted;
3637 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3648 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3659bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3660 SPIRVTypeInst ResType,
3661 MachineInstr &
I)
const {
3663 assert(
I.getNumOperands() == 3);
3665 auto Op =
I.getOperand(2);
3675 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3677 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3678 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3699 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3703 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3710bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3711 SPIRVTypeInst ResType,
3713 bool IsUnsigned)
const {
3714 return selectWaveReduce(
3715 ResVReg, ResType,
I, IsUnsigned,
3716 [&](
Register InputRegister,
bool IsUnsigned) {
3717 const bool IsFloatTy =
3719 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3720 : SPIRV::OpGroupNonUniformSMax;
3721 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3725bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3726 SPIRVTypeInst ResType,
3728 bool IsUnsigned)
const {
3729 return selectWaveReduce(
3730 ResVReg, ResType,
I, IsUnsigned,
3731 [&](
Register InputRegister,
bool IsUnsigned) {
3732 const bool IsFloatTy =
3734 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3735 : SPIRV::OpGroupNonUniformSMin;
3736 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3740bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3741 SPIRVTypeInst ResType,
3742 MachineInstr &
I)
const {
3743 return selectWaveReduce(ResVReg, ResType,
I,
false,
3744 [&](
Register InputRegister,
bool IsUnsigned) {
3746 InputRegister, SPIRV::OpTypeFloat);
3747 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3748 : SPIRV::OpGroupNonUniformIAdd;
3752bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3753 SPIRVTypeInst ResType,
3754 MachineInstr &
I)
const {
3755 return selectWaveReduce(ResVReg, ResType,
I,
false,
3756 [&](
Register InputRegister,
bool IsUnsigned) {
3758 InputRegister, SPIRV::OpTypeFloat);
3759 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3760 : SPIRV::OpGroupNonUniformIMul;
3764template <
typename PickOpcodeFn>
3765bool SPIRVInstructionSelector::selectWaveReduce(
3766 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3767 PickOpcodeFn &&PickOpcode)
const {
3768 assert(
I.getNumOperands() == 3);
3769 assert(
I.getOperand(2).isReg());
3771 Register InputRegister =
I.getOperand(2).getReg();
3775 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3778 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3784 .
addImm(SPIRV::GroupOperation::Reduce)
3785 .
addUse(
I.getOperand(2).getReg())
3790bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3791 SPIRVTypeInst ResType,
3793 unsigned Opcode)
const {
3794 return selectWaveReduce(
3795 ResVReg, ResType,
I,
false,
3796 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3799bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3800 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3801 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3802 [&](
Register InputRegister,
bool IsUnsigned) {
3804 InputRegister, SPIRV::OpTypeFloat);
3806 ? SPIRV::OpGroupNonUniformFAdd
3807 : SPIRV::OpGroupNonUniformIAdd;
3811bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3812 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3813 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3814 [&](
Register InputRegister,
bool IsUnsigned) {
3816 InputRegister, SPIRV::OpTypeFloat);
3818 ? SPIRV::OpGroupNonUniformFMul
3819 : SPIRV::OpGroupNonUniformIMul;
3823template <
typename PickOpcodeFn>
3824bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3825 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3826 PickOpcodeFn &&PickOpcode)
const {
3827 assert(
I.getNumOperands() == 3);
3828 assert(
I.getOperand(2).isReg());
3830 Register InputRegister =
I.getOperand(2).getReg();
3834 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3837 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3843 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3844 .
addUse(
I.getOperand(2).getReg())
3849bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3850 SPIRVTypeInst ResType,
3853 assert(
I.getNumOperands() == 3);
3854 assert(
I.getOperand(2).isReg());
3856 Register InputRegister =
I.getOperand(2).getReg();
3862 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3873bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
3874 SPIRVTypeInst ResType,
3881 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3886 : SPIRV::OpUConvert;
3890 ShiftOp = SPIRV::OpShiftRightLogicalV;
3895 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3896 TII.get(SPIRV::OpConstantComposite))
3899 for (
unsigned It = 0; It <
N; ++It)
3903 ShiftConst = CompositeReg;
3908 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3913 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3918 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3923 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3926bool SPIRVInstructionSelector::handle64BitOverflow(
3928 unsigned int Opcode,
3935 "handle64BitOverflow should only be used for integer types");
3937 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
3939 MachineIRBuilder MIRBuilder(
I);
3941 SPIRVTypeInst I64x2Type =
3943 SPIRVTypeInst Vec2ResType =
3946 std::vector<Register> PartialRegs;
3948 unsigned CurrentComponent = 0;
3949 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3953 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3954 TII.get(SPIRV::OpVectorShuffle))
3959 .
addImm(CurrentComponent)
3960 .
addImm(CurrentComponent + 1);
3970 PartialRegs.push_back(SubVecReg);
3973 if (CurrentComponent != ComponentCount) {
3979 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
3980 SPIRV::OpVectorExtractDynamic))
3989 PartialRegs.push_back(FinalElemResReg);
3993 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
3994 SPIRV::OpCompositeConstruct);
3997bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
3998 SPIRVTypeInst ResType,
4002 if (ComponentCount > 2)
4003 return handle64BitOverflow(
4004 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4006 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4008 MachineIRBuilder MIRBuilder(
I);
4012 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4016 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4021 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4028 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4029 TII.get(SPIRV::OpVectorShuffle))
4034 for (
unsigned J = 0; J < ComponentCount; ++J) {
4041 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4044bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4045 SPIRVTypeInst ResType,
4049 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4057bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4058 SPIRVTypeInst ResType,
4059 MachineInstr &
I)
const {
4060 Register OpReg =
I.getOperand(1).getReg();
4069 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4071 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4073 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4075 return SPIRVInstructionSelector::diagnoseUnsupported(
4076 I,
"G_BITREVERSE only support 16,32,64 bits.");
4080 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4091 unsigned AndOp = SPIRV::OpBitwiseAndS;
4092 unsigned OrOp = SPIRV::OpBitwiseOrS;
4093 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4094 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4096 AndOp = SPIRV::OpBitwiseAndV;
4097 OrOp = SPIRV::OpBitwiseOrV;
4098 ShlOp = SPIRV::OpShiftLeftLogicalV;
4099 ShrOp = SPIRV::OpShiftRightLogicalV;
4105 const unsigned Shift) ->
Register {
4113 Register MaskReg = CreateConst(Mask);
4114 Register ShiftReg = CreateConst(Shift);
4121 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4122 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4123 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4124 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4125 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4133 uint64_t
Mask = ~0ull;
4134 while ((Shift >>= 1) > 0) {
4141 return BuildCOPY(ResVReg, Result,
I);
4144bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4145 SPIRVTypeInst ResType,
4146 MachineInstr &
I)
const {
4147 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4148 "G_FREEZE must define and use a register");
4149 Register OpReg =
I.getOperand(1).getReg();
4153 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4166 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4167 if (
Def->getOpcode() == TargetOpcode::COPY)
4170 switch (
Def->getOpcode()) {
4171 case SPIRV::ASSIGN_TYPE:
4172 if (MachineInstr *AssignToDef =
4174 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4175 Reg =
Def->getOperand(2).getReg();
4178 case SPIRV::OpUndef:
4179 Reg =
Def->getOperand(1).getReg();
4182 unsigned DestOpCode;
4184 DestOpCode = SPIRV::OpConstantNull;
4185 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4186 "static undef/poison lowered to OpConstantNull\n");
4188 DestOpCode = TargetOpcode::COPY;
4190 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4191 "skipped, lowered as a copy of the operand\n");
4193 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4194 .
addDef(
I.getOperand(0).getReg())
4202bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4203 SPIRVTypeInst ResType,
4204 MachineInstr &
I)
const {
4206 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4208 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4212 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4217 for (
unsigned i =
I.getNumExplicitDefs();
4218 i <
I.getNumExplicitOperands() && IsConst; ++i)
4222 if (!IsConst &&
N < 2)
4223 return diagnoseUnsupported(
4224 I,
"There must be at least two constituent operands in a vector");
4229 for (
unsigned i =
I.getNumExplicitDefs();
4230 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4231 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4236 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4243 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4244 TII.get(IsConst ? SPIRV::OpConstantComposite
4245 : SPIRV::OpCompositeConstruct))
4248 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4249 MIB.
addUse(
I.getOperand(i).getReg());
4254bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4255 SPIRVTypeInst ResType,
4256 MachineInstr &
I)
const {
4258 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4260 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4266 if (!
I.getOperand(
OpIdx).isReg())
4273 if (!IsConst &&
N < 2)
4274 return diagnoseUnsupported(
4275 I,
"There must be at least two constituent operands in a vector");
4278 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4279 TII.get(IsConst ? SPIRV::OpConstantComposite
4280 : SPIRV::OpCompositeConstruct))
4283 for (
unsigned i = 0; i <
N; ++i)
4289bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4290 SPIRVTypeInst ResType,
4291 MachineInstr &
I)
const {
4295 if (ResType->
getOpcode() != SPIRV::OpTypeVector)
4297 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4299 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4300 TII.get(SPIRV::OpCompositeConstruct))
4310bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4311 SPIRVTypeInst ResType,
4312 MachineInstr &
I)
const {
4317 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4319 Opcode = SPIRV::OpDemoteToHelperInvocation;
4321 Opcode = SPIRV::OpKill;
4323 if (MachineInstr *NextI =
I.getNextNode()) {
4325 NextI->eraseFromParent();
4335bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4336 SPIRVTypeInst ResType,
unsigned CmpOpc,
4337 MachineInstr &
I)
const {
4338 Register Cmp0 =
I.getOperand(2).getReg();
4339 Register Cmp1 =
I.getOperand(3).getReg();
4342 "CMP operands should have the same type");
4343 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4353bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4354 SPIRVTypeInst ResType,
4355 MachineInstr &
I)
const {
4356 auto Pred =
I.getOperand(1).getPredicate();
4359 Register CmpOperand =
I.getOperand(2).getReg();
4364 Register Op1 =
I.getOperand(3).getReg();
4368 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4373 I.getOperand(3).setReg(NewOp1);
4379 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4383SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4384 SPIRVTypeInst ResType)
const {
4386 SPIRVTypeInst SpvI32Ty =
4389 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4396 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4399 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4402 .
addImm(APInt(32, Val).getZExtValue());
4404 GR.
add(ConstInt,
MI);
4411Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4412 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4414 SPIRVTypeInst SpvI32Ty =
4416 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4421 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4422 MachineInstr *
MI =
nullptr;
4426 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4430 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4431 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4437 GR.
add(ConstInt,
MI);
4442bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4443 SPIRVTypeInst ResType,
4444 MachineInstr &
I)
const {
4446 return selectCmp(ResVReg, ResType, CmpOp,
I);
4449bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4450 SPIRVTypeInst ResType,
4451 MachineInstr &
I)
const {
4453 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4460 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4461 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4464 MachineIRBuilder MIRBuilder(
I);
4471 APFloat ConstVal(3.3219280948873623);
4475 APFloat::rmNearestTiesToEven, &LosesInfo);
4479 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4480 ? SPIRV::OpVectorTimesScalar
4483 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4484 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4486 if (!selectExtInst(ResVReg, ResType,
I,
4487 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4497Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4498 MachineInstr &
I)
const {
4501 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4506bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4512 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4520 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4523 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4524 Def->getOpcode() == SPIRV::OpConstantI)
4537 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4538 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4540 Intrinsic::spv_const_composite)) {
4541 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4542 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4543 if (!IsZero(
Def->getOperand(i).getReg()))
4552Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4553 MachineInstr &
I)
const {
4557 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4562Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4563 MachineInstr &
I)
const {
4567 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4573 SPIRVTypeInst ResType,
4574 MachineInstr &
I)
const {
4578 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4583bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4584 SPIRVTypeInst ResType,
4585 MachineInstr &
I)
const {
4586 Register SelectFirstArg =
I.getOperand(2).getReg();
4587 Register SelectSecondArg =
I.getOperand(3).getReg();
4596 SPIRV::OpTypeVector;
4603 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4604 }
else if (IsPtrTy) {
4605 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4607 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4610 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4611 "boolean condition");
4613 Opcode = SPIRV::OpSelectSFSCond;
4614 }
else if (IsPtrTy) {
4615 Opcode = SPIRV::OpSelectSPSCond;
4617 Opcode = SPIRV::OpSelectSISCond;
4620 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4623 .
addUse(
I.getOperand(1).getReg())
4632bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4633 SPIRVTypeInst ResType,
4635 MachineInstr &InsertAt,
4636 bool IsSigned)
const {
4638 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4639 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4640 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4642 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4654bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4655 SPIRVTypeInst ResType,
4656 MachineInstr &
I,
bool IsSigned,
4657 unsigned Opcode)
const {
4658 Register SrcReg =
I.getOperand(1).getReg();
4664 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4669 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4671 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4674bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4675 SPIRVTypeInst ResType, MachineInstr &
I,
4676 bool IsSigned)
const {
4677 Register SrcReg =
I.getOperand(1).getReg();
4679 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4683 if (ResType == SrcType)
4684 return BuildCOPY(ResVReg, SrcReg,
I);
4686 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4687 return selectUnOp(ResVReg, ResType,
I, Opcode);
4690bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4691 SPIRVTypeInst ResType,
4693 bool IsSigned)
const {
4694 MachineIRBuilder MIRBuilder(
I);
4695 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4707 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4710 .
addUse(
I.getOperand(1).getReg())
4711 .
addUse(
I.getOperand(2).getReg())
4716 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4719 .
addUse(
I.getOperand(1).getReg())
4720 .
addUse(
I.getOperand(2).getReg())
4728 unsigned SelectOpcode =
4729 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4734 .
addUse(buildOnesVal(
true, ResType,
I))
4735 .
addUse(buildZerosVal(ResType,
I))
4742 .
addUse(buildOnesVal(
false, ResType,
I))
4747bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4750 SPIRVTypeInst IntTy,
4751 SPIRVTypeInst BoolTy)
const {
4754 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4755 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4757 Register One = buildOnesVal(
false, IntTy,
I);
4765 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4774bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4775 SPIRVTypeInst ResType,
4776 MachineInstr &
I)
const {
4777 Register IntReg =
I.getOperand(1).getReg();
4780 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4781 if (ArgType == ResType)
4782 return BuildCOPY(ResVReg, IntReg,
I);
4784 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4785 return selectUnOp(ResVReg, ResType,
I, Opcode);
4788bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4789 SPIRVTypeInst ResType,
4790 MachineInstr &
I)
const {
4791 unsigned Opcode =
I.getOpcode();
4792 unsigned TpOpcode = ResType->
getOpcode();
4794 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4795 assert(Opcode == TargetOpcode::G_CONSTANT &&
4796 I.getOperand(1).getCImm()->isZero());
4797 MachineBasicBlock &DepMBB =
I.getMF()->front();
4800 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4807 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4810bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4811 SPIRVTypeInst ResType,
4812 MachineInstr &
I)
const {
4813 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4820bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4821 SPIRVTypeInst ResType,
4822 MachineInstr &
I)
const {
4824 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4828 .
addUse(
I.getOperand(3).getReg())
4830 .
addUse(
I.getOperand(2).getReg());
4831 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4837bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4838 SPIRVTypeInst ResType,
4839 MachineInstr &
I)
const {
4840 Type *MaybeResTy =
nullptr;
4845 "Expected aggregate type for extractv instruction");
4847 SPIRV::AccessQualifier::ReadWrite,
false);
4851 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4854 .
addUse(
I.getOperand(2).getReg());
4855 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4861bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4862 SPIRVTypeInst ResType,
4863 MachineInstr &
I)
const {
4864 if (
getImm(
I.getOperand(4), MRI))
4865 return selectInsertVal(ResVReg, ResType,
I);
4867 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4870 .
addUse(
I.getOperand(2).getReg())
4871 .
addUse(
I.getOperand(3).getReg())
4872 .
addUse(
I.getOperand(4).getReg())
4877bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4878 SPIRVTypeInst ResType,
4879 MachineInstr &
I)
const {
4880 if (
getImm(
I.getOperand(3), MRI))
4881 return selectExtractVal(ResVReg, ResType,
I);
4883 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4886 .
addUse(
I.getOperand(2).getReg())
4887 .
addUse(
I.getOperand(3).getReg())
4892bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4893 SPIRVTypeInst ResType,
4894 MachineInstr &
I)
const {
4895 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4901 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4902 : SPIRV::OpAccessChain)
4903 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4904 :
SPIRV::OpPtrAccessChain);
4906 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4910 .
addUse(
I.getOperand(3).getReg());
4912 (Opcode == SPIRV::OpPtrAccessChain ||
4913 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4914 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4915 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4918 const unsigned StartingIndex =
4919 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4922 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4923 Res.addUse(
I.getOperand(i).getReg());
4924 Res.constrainAllUses(
TII,
TRI, RBI);
4929bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4931 unsigned Lim =
I.getNumExplicitOperands();
4932 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4933 Register OpReg =
I.getOperand(i).getReg();
4934 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4936 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4937 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4938 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4945 MachineFunction *MF =
I.getMF();
4957 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4958 TII.get(SPIRV::OpSpecConstantOp))
4961 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4963 GR.
add(OpDefine, MIB);
4969bool SPIRVInstructionSelector::selectDerivativeInst(
4970 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4971 const unsigned DPdOpCode)
const {
4974 if (!errorIfInstrOutsideShader(
I))
4980 Register SrcReg =
I.getOperand(2).getReg();
4985 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4988 .
addUse(
I.getOperand(2).getReg());
4990 MachineIRBuilder MIRBuilder(
I);
4993 if (componentCount != 1)
5001 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5006 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5011 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5019bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5020 SPIRVTypeInst ResType,
5021 MachineInstr &
I)
const {
5025 case Intrinsic::spv_load:
5026 return selectLoad(ResVReg, ResType,
I);
5027 case Intrinsic::spv_atomic_load:
5028 return selectAtomicLoad(ResVReg, ResType,
I);
5029 case Intrinsic::spv_store:
5030 return selectStore(
I);
5031 case Intrinsic::spv_atomic_store:
5032 return selectAtomicStore(
I);
5033 case Intrinsic::spv_extractv:
5034 return selectExtractVal(ResVReg, ResType,
I);
5035 case Intrinsic::spv_insertv:
5036 return selectInsertVal(ResVReg, ResType,
I);
5037 case Intrinsic::spv_extractelt:
5038 return selectExtractElt(ResVReg, ResType,
I);
5039 case Intrinsic::spv_insertelt:
5040 return selectInsertElt(ResVReg, ResType,
I);
5041 case Intrinsic::spv_gep:
5042 return selectGEP(ResVReg, ResType,
I);
5043 case Intrinsic::spv_bitcast: {
5044 Register OpReg =
I.getOperand(2).getReg();
5045 SPIRVTypeInst OpType =
5049 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5051 case Intrinsic::spv_unref_global:
5052 case Intrinsic::spv_init_global: {
5053 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
5058 Register GVarVReg =
MI->getOperand(0).getReg();
5059 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5064 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5066 MI->eraseFromParent();
5070 case Intrinsic::spv_undef: {
5071 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5077 case Intrinsic::spv_poison:
5078 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5083 case Intrinsic::spv_freeze:
5084 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5087 .
addUse(
I.getOperand(2).getReg())
5090 case Intrinsic::spv_named_boolean_spec_constant: {
5091 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5092 : SPIRV::OpSpecConstantFalse;
5094 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5095 .
addDef(
I.getOperand(0).getReg())
5098 unsigned SpecId =
I.getOperand(2).getImm();
5100 SPIRV::Decoration::SpecId, {SpecId});
5104 case Intrinsic::spv_const_composite: {
5106 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5112 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5114 std::function<bool(
Register)> HasSpecConstOperand =
5124 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5125 J < Def->getNumExplicitOperands(); ++J) {
5126 if (
Def->getOperand(J).isReg() &&
5127 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5133 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5134 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5135 : SPIRV::OpConstantComposite;
5136 unsigned ContinuedOpc = HasSpecConst
5137 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5138 : SPIRV::OpConstantCompositeContinuedINTEL;
5139 MachineIRBuilder MIR(
I);
5141 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5143 for (
auto *Instr : Instructions) {
5144 Instr->setDebugLoc(
I.getDebugLoc());
5149 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5156 case Intrinsic::spv_assign_name: {
5157 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5158 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5159 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5160 i <
I.getNumExplicitOperands(); ++i) {
5161 MIB.
addImm(
I.getOperand(i).getImm());
5166 case Intrinsic::spv_switch: {
5167 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5168 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5169 if (
I.getOperand(i).isReg())
5170 MIB.
addReg(
I.getOperand(i).getReg());
5171 else if (
I.getOperand(i).isCImm())
5172 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5173 else if (
I.getOperand(i).isMBB())
5174 MIB.
addMBB(
I.getOperand(i).getMBB());
5181 case Intrinsic::spv_loop_merge: {
5182 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5183 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5184 if (
I.getOperand(i).isMBB())
5185 MIB.
addMBB(
I.getOperand(i).getMBB());
5192 case Intrinsic::spv_loop_control_intel: {
5194 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5195 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5200 case Intrinsic::spv_selection_merge: {
5202 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5203 assert(
I.getOperand(1).isMBB() &&
5204 "operand 1 to spv_selection_merge must be a basic block");
5205 MIB.
addMBB(
I.getOperand(1).getMBB());
5206 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5210 case Intrinsic::spv_cmpxchg:
5211 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5212 case Intrinsic::spv_unreachable:
5213 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5216 case Intrinsic::spv_abort:
5217 return selectAbort(
I);
5218 case Intrinsic::spv_alloca:
5219 return selectFrameIndex(ResVReg, ResType,
I);
5220 case Intrinsic::spv_alloca_array:
5221 return selectAllocaArray(ResVReg, ResType,
I);
5222 case Intrinsic::spv_assume:
5224 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5225 .
addUse(
I.getOperand(1).getReg())
5230 case Intrinsic::spv_expect:
5232 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5235 .
addUse(
I.getOperand(2).getReg())
5236 .
addUse(
I.getOperand(3).getReg())
5241 case Intrinsic::arithmetic_fence:
5242 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5243 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5246 .
addUse(
I.getOperand(2).getReg())
5250 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5252 case Intrinsic::spv_thread_id:
5258 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5260 case Intrinsic::spv_thread_id_in_group:
5266 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5268 case Intrinsic::spv_group_id:
5274 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5276 case Intrinsic::spv_flattened_thread_id_in_group:
5283 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5285 case Intrinsic::spv_workgroup_size:
5286 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5288 case Intrinsic::spv_global_size:
5289 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5291 case Intrinsic::spv_global_offset:
5292 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5294 case Intrinsic::spv_num_workgroups:
5295 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5297 case Intrinsic::spv_subgroup_size:
5298 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5300 case Intrinsic::spv_num_subgroups:
5301 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5303 case Intrinsic::spv_subgroup_id:
5304 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5305 case Intrinsic::spv_subgroup_local_invocation_id:
5306 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5307 ResVReg, ResType,
I);
5308 case Intrinsic::spv_subgroup_max_size:
5309 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5311 case Intrinsic::spv_fdot:
5312 return selectFloatDot(ResVReg, ResType,
I);
5313 case Intrinsic::spv_udot:
5314 case Intrinsic::spv_sdot:
5315 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5317 return selectIntegerDot(ResVReg, ResType,
I,
5318 IID == Intrinsic::spv_sdot);
5319 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5320 case Intrinsic::spv_dot4add_i8packed:
5321 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5323 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5324 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5325 case Intrinsic::spv_dot4add_u8packed:
5326 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5328 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5329 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5330 case Intrinsic::spv_all:
5331 return selectAll(ResVReg, ResType,
I);
5332 case Intrinsic::spv_any:
5333 return selectAny(ResVReg, ResType,
I);
5334 case Intrinsic::spv_cross:
5335 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
5336 case Intrinsic::spv_distance:
5337 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5338 case Intrinsic::spv_lerp:
5339 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5340 case Intrinsic::spv_length:
5341 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5342 case Intrinsic::spv_degrees:
5343 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5344 case Intrinsic::spv_faceforward:
5345 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5346 case Intrinsic::spv_frac:
5347 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5348 case Intrinsic::spv_isinf:
5349 return selectOpIsInf(ResVReg, ResType,
I);
5350 case Intrinsic::spv_isnan:
5351 return selectOpIsNan(ResVReg, ResType,
I);
5352 case Intrinsic::spv_isfinite:
5353 return selectOpIsFinite(ResVReg, ResType,
I);
5354 case Intrinsic::spv_isnormal:
5355 return selectOpIsNormal(ResVReg, ResType,
I);
5356 case Intrinsic::spv_normalize:
5357 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5358 case Intrinsic::spv_refract:
5359 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5360 case Intrinsic::spv_reflect:
5361 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5362 case Intrinsic::spv_rsqrt:
5363 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5364 case Intrinsic::spv_sign:
5365 return selectSign(ResVReg, ResType,
I);
5366 case Intrinsic::spv_smoothstep:
5367 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5368 case Intrinsic::spv_firstbituhigh:
5369 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5370 case Intrinsic::spv_firstbitshigh:
5371 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5372 case Intrinsic::spv_firstbitlow:
5373 return selectFirstBitLow(ResVReg, ResType,
I);
5374 case Intrinsic::spv_all_memory_barrier:
5375 return selectBarrierInst(
I, SPIRV::Scope::Device,
5376 SPIRV::MemorySemantics::UniformMemory |
5377 SPIRV::MemorySemantics::ImageMemory |
5378 SPIRV::MemorySemantics::WorkgroupMemory,
5380 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5381 return selectBarrierInst(
I, SPIRV::Scope::Device,
5382 SPIRV::MemorySemantics::UniformMemory |
5383 SPIRV::MemorySemantics::ImageMemory |
5384 SPIRV::MemorySemantics::WorkgroupMemory,
5386 case Intrinsic::spv_device_memory_barrier:
5387 return selectBarrierInst(
I, SPIRV::Scope::Device,
5388 SPIRV::MemorySemantics::UniformMemory |
5389 SPIRV::MemorySemantics::ImageMemory,
5391 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5392 return selectBarrierInst(
I, SPIRV::Scope::Device,
5393 SPIRV::MemorySemantics::UniformMemory |
5394 SPIRV::MemorySemantics::ImageMemory,
5396 case Intrinsic::spv_group_memory_barrier:
5397 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5398 SPIRV::MemorySemantics::WorkgroupMemory,
5400 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5401 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5402 SPIRV::MemorySemantics::WorkgroupMemory,
5404 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5405 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5406 SPIRV::StorageClass::StorageClass ResSC =
5409 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5410 "from the Generic storage class");
5411 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5419 case Intrinsic::spv_lifetime_start:
5420 case Intrinsic::spv_lifetime_end: {
5421 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5422 : SPIRV::OpLifetimeStop;
5423 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5424 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5433 case Intrinsic::spv_saturate:
5434 return selectSaturate(ResVReg, ResType,
I);
5435 case Intrinsic::spv_nclamp:
5436 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5437 case Intrinsic::spv_uclamp:
5438 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5439 case Intrinsic::spv_sclamp:
5440 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5441 case Intrinsic::spv_subgroup_prefix_bit_count:
5442 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5443 case Intrinsic::spv_wave_active_countbits:
5444 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5445 case Intrinsic::spv_wave_all_equal:
5446 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5447 case Intrinsic::spv_wave_all:
5448 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5449 case Intrinsic::spv_wave_any:
5450 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5451 case Intrinsic::spv_subgroup_ballot:
5452 return selectWaveOpInst(ResVReg, ResType,
I,
5453 SPIRV::OpGroupNonUniformBallot);
5454 case Intrinsic::spv_wave_is_first_lane:
5455 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5456 case Intrinsic::spv_wave_reduce_or:
5457 return selectWaveReduceOp(ResVReg, ResType,
I,
5458 SPIRV::OpGroupNonUniformBitwiseOr);
5459 case Intrinsic::spv_wave_reduce_xor:
5460 return selectWaveReduceOp(ResVReg, ResType,
I,
5461 SPIRV::OpGroupNonUniformBitwiseXor);
5462 case Intrinsic::spv_wave_reduce_and:
5463 return selectWaveReduceOp(ResVReg, ResType,
I,
5464 SPIRV::OpGroupNonUniformBitwiseAnd);
5465 case Intrinsic::spv_interlocked_add:
5466 return selectInterlockedOp(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
5467 case Intrinsic::spv_interlocked_or:
5468 return selectInterlockedOp(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
5469 case Intrinsic::spv_wave_reduce_umax:
5470 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5471 case Intrinsic::spv_wave_reduce_max:
5472 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5473 case Intrinsic::spv_wave_reduce_umin:
5474 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5475 case Intrinsic::spv_wave_reduce_min:
5476 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5477 case Intrinsic::spv_wave_reduce_sum:
5478 return selectWaveReduceSum(ResVReg, ResType,
I);
5479 case Intrinsic::spv_wave_product:
5480 return selectWaveReduceProduct(ResVReg, ResType,
I);
5481 case Intrinsic::spv_wave_readlane:
5482 return selectWaveOpInst(ResVReg, ResType,
I,
5483 SPIRV::OpGroupNonUniformShuffle);
5484 case Intrinsic::spv_wave_prefix_sum:
5485 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5486 case Intrinsic::spv_wave_prefix_product:
5487 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5488 case Intrinsic::spv_quad_read_across_x: {
5489 return selectQuadSwap(ResVReg, ResType,
I, 0);
5491 case Intrinsic::spv_quad_read_across_y: {
5492 return selectQuadSwap(ResVReg, ResType,
I, 1);
5494 case Intrinsic::spv_quad_read_across_diagonal: {
5495 return selectQuadSwap(ResVReg, ResType,
I, 2);
5497 case Intrinsic::spv_step:
5498 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5499 case Intrinsic::spv_radians:
5500 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5504 case Intrinsic::instrprof_increment:
5505 case Intrinsic::instrprof_increment_step:
5506 case Intrinsic::instrprof_value_profile:
5509 case Intrinsic::spv_value_md:
5511 case Intrinsic::spv_resource_handlefrombinding: {
5512 return selectHandleFromBinding(ResVReg, ResType,
I);
5514 case Intrinsic::spv_resource_counterhandlefrombinding:
5515 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5516 case Intrinsic::spv_resource_updatecounter:
5517 return selectUpdateCounter(ResVReg, ResType,
I);
5518 case Intrinsic::spv_resource_store_typedbuffer: {
5519 return selectImageWriteIntrinsic(
I);
5521 case Intrinsic::spv_resource_load_typedbuffer: {
5522 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5524 case Intrinsic::spv_resource_load_level: {
5525 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5527 case Intrinsic::spv_resource_getdimensions_x:
5528 case Intrinsic::spv_resource_getdimensions_xy:
5529 case Intrinsic::spv_resource_getdimensions_xyz: {
5530 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5532 case Intrinsic::spv_resource_getdimensions_levels_x:
5533 case Intrinsic::spv_resource_getdimensions_levels_xy:
5534 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5535 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5537 case Intrinsic::spv_resource_getdimensions_ms_xy:
5538 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5539 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5541 case Intrinsic::spv_resource_calculate_lod:
5542 case Intrinsic::spv_resource_calculate_lod_unclamped:
5543 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5544 case Intrinsic::spv_resource_sample:
5545 case Intrinsic::spv_resource_sample_clamp:
5546 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5547 case Intrinsic::spv_resource_samplebias:
5548 case Intrinsic::spv_resource_samplebias_clamp:
5549 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5550 case Intrinsic::spv_resource_samplegrad:
5551 case Intrinsic::spv_resource_samplegrad_clamp:
5552 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5553 case Intrinsic::spv_resource_samplelevel:
5554 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5555 case Intrinsic::spv_resource_samplecmp:
5556 case Intrinsic::spv_resource_samplecmp_clamp:
5557 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5558 case Intrinsic::spv_resource_samplecmplevelzero:
5559 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5560 case Intrinsic::spv_resource_gather:
5561 case Intrinsic::spv_resource_gather_cmp:
5562 return selectGatherIntrinsic(ResVReg, ResType,
I);
5563 case Intrinsic::spv_resource_getbasepointer:
5564 case Intrinsic::spv_resource_getpointer: {
5565 return selectResourceGetPointer(ResVReg, ResType,
I);
5567 case Intrinsic::spv_pushconstant_getpointer: {
5568 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5570 case Intrinsic::spv_discard: {
5571 return selectDiscard(ResVReg, ResType,
I);
5573 case Intrinsic::spv_resource_nonuniformindex: {
5574 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5576 case Intrinsic::spv_unpackhalf2x16: {
5577 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5579 case Intrinsic::spv_packhalf2x16: {
5580 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5582 case Intrinsic::spv_ddx:
5583 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5584 case Intrinsic::spv_ddy:
5585 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5586 case Intrinsic::spv_ddx_coarse:
5587 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5588 case Intrinsic::spv_ddy_coarse:
5589 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5590 case Intrinsic::spv_ddx_fine:
5591 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5592 case Intrinsic::spv_ddy_fine:
5593 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5594 case Intrinsic::spv_fwidth:
5595 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5596 case Intrinsic::spv_masked_gather:
5597 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5598 return selectMaskedGather(ResVReg, ResType,
I);
5599 return diagnoseUnsupported(
5600 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5601 case Intrinsic::spv_masked_scatter:
5602 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5603 return selectMaskedScatter(
I);
5604 return diagnoseUnsupported(
5605 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5606 case Intrinsic::returnaddress:
5607 case Intrinsic::frameaddress: {
5609 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5616 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5621bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5622 SPIRVTypeInst ResType,
5623 MachineInstr &
I)
const {
5626 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5633bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5634 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5636 assert(Intr.getIntrinsicID() ==
5637 Intrinsic::spv_resource_counterhandlefrombinding);
5640 Register MainHandleReg = Intr.getOperand(2).getReg();
5642 assert(MainHandleDef->getIntrinsicID() ==
5643 Intrinsic::spv_resource_handlefrombinding);
5647 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5648 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5649 std::string CounterName =
5654 MachineIRBuilder MIRBuilder(
I);
5656 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5658 ArraySize, IndexReg, CounterName, MIRBuilder);
5660 return BuildCOPY(ResVReg, CounterVarReg,
I);
5663bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5664 SPIRVTypeInst ResType,
5665 MachineInstr &
I)
const {
5667 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5669 Register CounterHandleReg = Intr.getOperand(2).getReg();
5670 Register IncrReg = Intr.getOperand(3).getReg();
5677 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5678 assert(CounterVarPointeeType &&
5679 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5680 "Counter variable must be a struct");
5682 SPIRV::StorageClass::StorageBuffer &&
5683 "Counter variable must be in the storage buffer storage class");
5685 "Counter variable must have exactly 1 member in the struct");
5686 const SPIRVTypeInst MemberType =
5689 "Counter variable struct must have a single i32 member");
5693 MachineIRBuilder MIRBuilder(
I);
5695 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5698 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5704 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5707 .
addUse(CounterHandleReg)
5714 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5717 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5720 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5729 return BuildCOPY(ResVReg, AtomicRes,
I);
5737 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5745bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5746 SPIRVTypeInst ResType,
5747 MachineInstr &
I)
const {
5755 Register ImageReg =
I.getOperand(2).getReg();
5763 Register IdxReg =
I.getOperand(3).getReg();
5765 MachineInstr &Pos =
I;
5767 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5771bool SPIRVInstructionSelector::generateSampleImage(
5774 DebugLoc Loc, MachineInstr &Pos)
const {
5785 if (!loadHandleBeforePosition(NewSamplerReg,
5791 MachineIRBuilder MIRBuilder(Pos);
5804 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5805 ImOps.Lod.has_value();
5806 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5807 : SPIRV::OpImageSampleImplicitLod;
5809 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5810 : SPIRV::OpImageSampleDrefImplicitLod;
5819 MIB.
addUse(*ImOps.Compare);
5821 uint32_t ImageOperands = 0;
5823 ImageOperands |= SPIRV::ImageOperand::Bias;
5825 ImageOperands |= SPIRV::ImageOperand::Lod;
5826 if (ImOps.GradX && ImOps.GradY)
5827 ImageOperands |= SPIRV::ImageOperand::Grad;
5828 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5830 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5833 "Non-constant offsets are not supported in sample instructions.");
5838 ImageOperands |= SPIRV::ImageOperand::MinLod;
5840 if (ImageOperands != 0) {
5841 MIB.
addImm(ImageOperands);
5842 if (ImageOperands & SPIRV::ImageOperand::Bias)
5844 if (ImageOperands & SPIRV::ImageOperand::Lod)
5846 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5847 MIB.
addUse(*ImOps.GradX);
5848 MIB.
addUse(*ImOps.GradY);
5851 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5852 MIB.
addUse(*ImOps.Offset);
5853 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5854 MIB.
addUse(*ImOps.MinLod);
5861bool SPIRVInstructionSelector::selectImageQuerySize(
5863 std::optional<Register> LodReg)
const {
5865 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5868 "ImageReg is not an image type.");
5870 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5872 unsigned NumComponents = 0;
5874 case SPIRV::Dim::DIM_1D:
5875 case SPIRV::Dim::DIM_Buffer:
5876 NumComponents =
IsArray ? 2 : 1;
5878 case SPIRV::Dim::DIM_2D:
5879 case SPIRV::Dim::DIM_Cube:
5880 case SPIRV::Dim::DIM_Rect:
5881 NumComponents =
IsArray ? 3 : 2;
5883 case SPIRV::Dim::DIM_3D:
5887 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5892 SPIRVTypeInst ResType =
5897 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5907bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5908 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5909 Register ImageReg =
I.getOperand(2).getReg();
5916 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5919bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5920 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5921 Register ImageReg =
I.getOperand(2).getReg();
5930 Register LodReg =
I.getOperand(3).getReg();
5933 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5935 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5942 TII.get(SPIRV::OpImageQueryLevels))
5949 TII.get(SPIRV::OpCompositeConstruct))
5959bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5960 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5961 Register ImageReg =
I.getOperand(2).getReg();
5972 "OpImageQuerySamples requires a multisampled image");
5974 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5982 TII.get(SPIRV::OpImageQuerySamples))
5989 TII.get(SPIRV::OpCompositeConstruct))
5999bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6000 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6001 Register ImageReg =
I.getOperand(2).getReg();
6002 Register SamplerReg =
I.getOperand(3).getReg();
6003 Register CoordinateReg =
I.getOperand(4).getReg();
6019 if (!loadHandleBeforePosition(
6024 MachineIRBuilder MIRBuilder(
I);
6030 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6040 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6047 unsigned ExtractedIndex =
6049 Intrinsic::spv_resource_calculate_lod_unclamped
6053 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6054 TII.get(SPIRV::OpCompositeExtract))
6064bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6065 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6066 Register ImageReg =
I.getOperand(2).getReg();
6067 Register SamplerReg =
I.getOperand(3).getReg();
6068 Register CoordinateReg =
I.getOperand(4).getReg();
6069 ImageOperands ImOps;
6070 if (
I.getNumOperands() > 5)
6071 ImOps.Offset =
I.getOperand(5).getReg();
6072 if (
I.getNumOperands() > 6)
6073 ImOps.MinLod =
I.getOperand(6).getReg();
6074 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6075 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6078bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6079 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6080 Register ImageReg =
I.getOperand(2).getReg();
6081 Register SamplerReg =
I.getOperand(3).getReg();
6082 Register CoordinateReg =
I.getOperand(4).getReg();
6083 ImageOperands ImOps;
6084 ImOps.Bias =
I.getOperand(5).getReg();
6085 if (
I.getNumOperands() > 6)
6086 ImOps.Offset =
I.getOperand(6).getReg();
6087 if (
I.getNumOperands() > 7)
6088 ImOps.MinLod =
I.getOperand(7).getReg();
6089 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6090 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6093bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6094 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6095 Register ImageReg =
I.getOperand(2).getReg();
6096 Register SamplerReg =
I.getOperand(3).getReg();
6097 Register CoordinateReg =
I.getOperand(4).getReg();
6098 ImageOperands ImOps;
6099 ImOps.GradX =
I.getOperand(5).getReg();
6100 ImOps.GradY =
I.getOperand(6).getReg();
6101 if (
I.getNumOperands() > 7)
6102 ImOps.Offset =
I.getOperand(7).getReg();
6103 if (
I.getNumOperands() > 8)
6104 ImOps.MinLod =
I.getOperand(8).getReg();
6105 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6106 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6109bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6110 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6111 Register ImageReg =
I.getOperand(2).getReg();
6112 Register SamplerReg =
I.getOperand(3).getReg();
6113 Register CoordinateReg =
I.getOperand(4).getReg();
6114 ImageOperands ImOps;
6115 ImOps.Lod =
I.getOperand(5).getReg();
6116 if (
I.getNumOperands() > 6)
6117 ImOps.Offset =
I.getOperand(6).getReg();
6118 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6119 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6122bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6123 SPIRVTypeInst ResType,
6124 MachineInstr &
I)
const {
6125 Register ImageReg =
I.getOperand(2).getReg();
6126 Register SamplerReg =
I.getOperand(3).getReg();
6127 Register CoordinateReg =
I.getOperand(4).getReg();
6128 ImageOperands ImOps;
6129 ImOps.Compare =
I.getOperand(5).getReg();
6130 if (
I.getNumOperands() > 6)
6131 ImOps.Offset =
I.getOperand(6).getReg();
6132 if (
I.getNumOperands() > 7)
6133 ImOps.MinLod =
I.getOperand(7).getReg();
6134 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6135 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6138bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6139 SPIRVTypeInst ResType,
6140 MachineInstr &
I)
const {
6141 Register ImageReg =
I.getOperand(2).getReg();
6142 Register CoordinateReg =
I.getOperand(3).getReg();
6143 Register LodReg =
I.getOperand(4).getReg();
6145 ImageOperands ImOps;
6147 if (
I.getNumOperands() > 5)
6148 ImOps.Offset =
I.getOperand(5).getReg();
6160 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6161 I.getDebugLoc(),
I, &ImOps);
6164bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6165 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6166 Register ImageReg =
I.getOperand(2).getReg();
6167 Register SamplerReg =
I.getOperand(3).getReg();
6168 Register CoordinateReg =
I.getOperand(4).getReg();
6169 ImageOperands ImOps;
6170 ImOps.Compare =
I.getOperand(5).getReg();
6171 if (
I.getNumOperands() > 6)
6172 ImOps.Offset =
I.getOperand(6).getReg();
6175 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6176 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6179bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6180 SPIRVTypeInst ResType,
6181 MachineInstr &
I)
const {
6182 Register ImageReg =
I.getOperand(2).getReg();
6183 Register SamplerReg =
I.getOperand(3).getReg();
6184 Register CoordinateReg =
I.getOperand(4).getReg();
6187 "ImageReg is not an image type.");
6192 ComponentOrCompareReg =
I.getOperand(5).getReg();
6193 OffsetReg =
I.getOperand(6).getReg();
6196 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6200 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6201 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6202 Dim != SPIRV::Dim::DIM_Rect) {
6204 "Gather operations are only supported for 2D, Cube, and Rect images.");
6211 if (!loadHandleBeforePosition(
6216 MachineIRBuilder MIRBuilder(
I);
6217 SPIRVTypeInst SampledImageType =
6222 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6230 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6232 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6234 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6239 .
addUse(ComponentOrCompareReg);
6241 uint32_t ImageOperands = 0;
6242 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6243 if (Dim == SPIRV::Dim::DIM_Cube) {
6245 "Gather operations with offset are not supported for Cube images.");
6249 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6251 ImageOperands |= SPIRV::ImageOperand::Offset;
6255 if (ImageOperands != 0) {
6256 MIB.
addImm(ImageOperands);
6258 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6266bool SPIRVInstructionSelector::generateImageReadOrFetch(
6269 const ImageOperands *ImOps)
const {
6272 "ImageReg is not an image type.");
6274 bool IsSignedInteger =
6279 bool IsFetch = (SampledOp.getImm() == 1);
6281 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6282 uint32_t ImageOperandsMask = 0;
6283 if (IsSignedInteger)
6284 ImageOperandsMask |= 0x1000;
6286 if (IsFetch && ImOps) {
6288 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6289 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6291 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6293 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6297 if (ImageOperandsMask != 0) {
6298 MIB.
addImm(ImageOperandsMask);
6299 if (IsFetch && ImOps) {
6302 if (ImOps->Offset &&
6303 (ImageOperandsMask &
6304 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6305 MIB.
addUse(*ImOps->Offset);
6314 SPIRVTypeInst SampledType =
6317 SPIRVTypeInst ReadType =
6318 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6319 bool ReadTypeMatchesResult = ReadType == ResType;
6321 Register ReadReg = ReadTypeMatchesResult
6327 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6333 BMI.constrainAllUses(
TII,
TRI, RBI);
6335 if (ReadTypeMatchesResult)
6348 if (ResultSize == 1) {
6357 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6360bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6361 SPIRVTypeInst ResType,
6362 MachineInstr &
I)
const {
6363 Register ResourcePtr =
I.getOperand(2).getReg();
6365 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6374 MachineIRBuilder MIRBuilder(
I);
6379 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6385 if (
I.getNumExplicitOperands() > 3) {
6386 Register IndexReg =
I.getOperand(3).getReg();
6393bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6394 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6399bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6400 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6401 Register ObjReg =
I.getOperand(2).getReg();
6402 if (!BuildCOPY(ResVReg, ObjReg,
I))
6412 decorateUsesAsNonUniform(ResVReg);
6416void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6419 while (WorkList.
size() > 0) {
6423 bool IsDecorated =
false;
6425 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6426 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6432 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6434 if (ResultReg == CurrentReg)
6442 SPIRV::Decoration::NonUniformEXT, {});
6447bool SPIRVInstructionSelector::extractSubvector(
6449 MachineInstr &InsertionPoint)
const {
6451 [[maybe_unused]] uint64_t InputSize =
6454 assert(InputSize > 1 &&
"The input must be a vector.");
6455 assert(ResultSize > 1 &&
"The result must be a vector.");
6456 assert(ResultSize < InputSize &&
6457 "Cannot extract more element than there are in the input.");
6461 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6464 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6473 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6475 TII.get(SPIRV::OpCompositeConstruct))
6479 for (
Register ComponentReg : ComponentRegisters)
6480 MIB.
addUse(ComponentReg);
6485bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6486 MachineInstr &
I)
const {
6493 Register ImageReg =
I.getOperand(1).getReg();
6501 Register CoordinateReg =
I.getOperand(2).getReg();
6502 Register DataReg =
I.getOperand(3).getReg();
6505 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6513Register SPIRVInstructionSelector::buildPointerToResource(
6514 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6515 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6516 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6518 if (ArraySize == 1) {
6519 SPIRVTypeInst PtrType =
6522 "SpirvResType did not have an explicit layout.");
6527 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6528 SPIRVTypeInst VarPointerType =
6531 VarPointerType, Set,
Binding, Name, MIRBuilder);
6533 SPIRVTypeInst ResPointerType =
6546bool SPIRVInstructionSelector::selectFirstBitSet16(
6547 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6548 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6550 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6554 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6557bool SPIRVInstructionSelector::selectFirstBitSet32(
6559 unsigned BitSetOpcode)
const {
6560 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6563 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6570bool SPIRVInstructionSelector::selectFirstBitSet64(
6572 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6585 if (ComponentCount > 2) {
6586 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6588 unsigned Opcode) ->
bool {
6589 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6593 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6597 MachineIRBuilder MIRBuilder(
I);
6599 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6603 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6609 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6616 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6619 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6620 SPIRV::OpVectorExtractDynamic))
6622 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6623 SPIRV::OpVectorExtractDynamic))
6627 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6628 TII.get(SPIRV::OpVectorShuffle))
6636 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6642 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6643 TII.get(SPIRV::OpVectorShuffle))
6651 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6671 SelectOp = SPIRV::OpSelectSISCond;
6672 AddOp = SPIRV::OpIAddS;
6680 SelectOp = SPIRV::OpSelectVIVCond;
6681 AddOp = SPIRV::OpIAddV;
6687 Register RegSecondaryOffset = Reg0;
6691 if (SwapPrimarySide) {
6692 PrimaryReg = LowReg;
6693 SecondaryReg = HighReg;
6694 RegPrimaryOffset = Reg0;
6695 RegSecondaryOffset = Reg32;
6700 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6701 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6706 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6707 SPIRV::OpINotEqual))
6714 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6715 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6720 if (SwapPrimarySide) {
6722 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6723 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6734 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6735 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6740 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6741 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6744 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6748bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6749 SPIRVTypeInst ResType,
6751 bool IsSigned)
const {
6753 Register OpReg =
I.getOperand(2).getReg();
6756 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6757 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6761 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6763 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6765 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6768 return diagnoseUnsupported(
6770 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6774bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6775 SPIRVTypeInst ResType,
6776 MachineInstr &
I)
const {
6778 Register OpReg =
I.getOperand(2).getReg();
6783 unsigned ExtendOpcode = SPIRV::OpUConvert;
6784 unsigned BitSetOpcode = GL::FindILsb;
6788 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6790 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6792 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6795 return diagnoseUnsupported(
I,
6796 "spv_firstbitlow only supports 16,32,64 bits.");
6800bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6801 SPIRVTypeInst ResType,
6802 MachineInstr &
I)
const {
6806 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6809 .
addUse(
I.getOperand(2).getReg())
6812 unsigned Alignment =
I.getOperand(3).getImm();
6826 while (!Worklist.
empty()) {
6828 switch (
T->getOpcode()) {
6829 case SPIRV::OpTypeInt:
6830 case SPIRV::OpTypeFloat:
6831 case SPIRV::OpTypePointer:
6833 case SPIRV::OpTypeVector:
6834 case SPIRV::OpTypeMatrix:
6835 case SPIRV::OpTypeArray: {
6836 Register OperandReg =
T->getOperand(1).getReg();
6840 case SPIRV::OpTypeStruct:
6841 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6842 Register OperandReg =
T->getOperand(Idx).getReg();
6854bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6855 assert(
I.getNumExplicitOperands() == 2);
6857 Register MsgReg =
I.getOperand(1).getReg();
6859 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6862 return diagnoseUnsupported(
6864 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6865 "scalar, pointer, vector, matrix, or aggregate of such types)");
6868 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6875bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6884 uint32_t MsgVal = ~0
u;
6885 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6886 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6889 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6892 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6899bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6900 SPIRVTypeInst ResType,
6901 MachineInstr &
I)
const {
6905 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6908 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6911 unsigned Alignment =
I.getOperand(2).getImm();
6918bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6923 const MachineInstr *PrevI =
I.getPrevNode();
6925 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6929 .
addMBB(
I.getOperand(0).getMBB())
6934 .
addMBB(
I.getOperand(0).getMBB())
6939bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6950 const MachineInstr *NextI =
I.getNextNode();
6952 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6958 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6960 .
addUse(
I.getOperand(0).getReg())
6961 .
addMBB(
I.getOperand(1).getMBB())
6967bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6968 MachineInstr &
I)
const {
6970 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6972 const unsigned NumOps =
I.getNumOperands();
6973 for (
unsigned i = 1; i <
NumOps; i += 2) {
6974 MIB.
addUse(
I.getOperand(i + 0).getReg());
6975 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6981bool SPIRVInstructionSelector::selectGlobalValue(
6982 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6984 MachineIRBuilder MIRBuilder(
I);
6985 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6988 std::string GlobalIdent;
6990 unsigned &
ID = UnnamedGlobalIDs[GV];
6992 ID = UnnamedGlobalIDs.
size();
6993 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
7019 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7026 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7031 MachineInstrBuilder MIB1 =
7032 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7035 MachineInstrBuilder MIB2 =
7037 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7041 GR.
add(ConstVal, MIB2);
7049 MachineInstrBuilder MIB3 =
7050 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7053 GR.
add(ConstVal, MIB3);
7059 assert(NewReg != ResVReg);
7060 return BuildCOPY(ResVReg, NewReg,
I);
7070 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7073 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7079 SPIRVTypeInst ResType =
7083 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7088 if (
GlobalVar->isExternallyInitialized() &&
7089 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7090 constexpr unsigned ReadWriteINTEL = 3u;
7093 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7099bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7100 SPIRVTypeInst ResType,
7101 MachineInstr &
I)
const {
7103 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7111 MachineIRBuilder MIRBuilder(
I);
7116 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7119 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7121 .
add(
I.getOperand(1))
7126 ResType->
getOpcode() == SPIRV::OpTypeFloat);
7136 APFloat::rmNearestTiesToEven, &LosesInfo);
7140 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
7141 ? SPIRV::OpVectorTimesScalar
7152bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7153 SPIRVTypeInst ResType,
7154 MachineInstr &
I)
const {
7157 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7163 Register ExpReg =
I.getOperand(2).getReg();
7165 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7166 SPIRV::OpConvertSToF))
7168 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7175bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7176 SPIRVTypeInst ResType,
7177 MachineInstr &
I)
const {
7193 MachineIRBuilder MIRBuilder(
I);
7194 SPIRVTypeInst FloatType =
7198 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7211 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7213 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
7216 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7222 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7225 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7228 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7232 Register IntegralPartReg =
I.getOperand(1).getReg();
7235 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7245 assert(
false &&
"GLSL::Modf is deprecated.");
7256bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7257 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7258 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7259 MachineIRBuilder MIRBuilder(
I);
7260 const SPIRVTypeInst Vec3Ty =
7263 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7275 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7279 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7285 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7292 assert(
I.getOperand(2).isReg());
7293 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7297 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7308bool SPIRVInstructionSelector::loadBuiltinInputID(
7309 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7310 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7311 MachineIRBuilder MIRBuilder(
I);
7313 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7328 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7332 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7341SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7342 MachineInstr &
I)
const {
7343 MachineIRBuilder MIRBuilder(
I);
7344 if (
Type->getOpcode() != SPIRV::OpTypeVector)
7354bool SPIRVInstructionSelector::loadHandleBeforePosition(
7355 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7356 MachineInstr &Pos)
const {
7359 Intrinsic::spv_resource_handlefrombinding);
7367 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7368 MachineIRBuilder MIRBuilder(HandleDef);
7369 SPIRVTypeInst VarType = ResType;
7370 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7372 if (IsStructuredBuffer) {
7377 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7379 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7382 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7383 ArraySize, IndexReg, Name, MIRBuilder);
7387 uint32_t LoadOpcode =
7388 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7398bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7399 MachineInstr &
I)
const {
7401 return diagnoseUnsupported(
7402 I,
"this instruction is only supported in shaders.");
7407InstructionSelector *
7411 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
MachineInstrBuilder & UseMI
#define GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
This file declares a class to represent arbitrary precision floating point values and provide a varie...
static bool selectUnmergeValues(MachineInstrBuilder &MIB, const ARMBaseInstrInfo &TII, MachineRegisterInfo &MRI, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
static uint8_t SwapBits(uint8_t Val)
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
DXIL Resource Implicit Binding
Declares convenience wrapper classes for interpreting MachineInstr instances as specific generic oper...
const HexagonInstrInfo * TII
LLVMTypeRef LLVMIntType(unsigned NumBits)
const size_t AbstractManglingParser< Derived, Alloc >::NumOps
Loop::LoopBounds::Direction Direction
Register const TargetRegisterInfo * TRI
Promote Memory to Register
MachineInstr unsigned OpIdx
uint64_t IntrinsicInst * II
static StringRef getName(Value *V)
static unsigned getFCmpOpcode(CmpInst::Predicate Pred, unsigned Size)
static bool isConcreteSPIRVType(SPIRVTypeInst Ty, const SPIRVGlobalRegistry &GR)
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)
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 ManagedStatic< cl::opt< FnT >, OptCreatorT > CallbackFunction
static const fltSemantics & IEEEsingle()
static const fltSemantics & BFloat()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
const fltSemantics & getSemantics() const
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.
Represent a constant reference to an array (0 or more elements consecutively in memory),...
BlockFrequencyInfo pass uses BlockFrequencyInfoImpl implementation to estimate IR basic block frequen...
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
Represents a call to an intrinsic.
Intrinsic::ID getIntrinsicID() const
unsigned getAddressSpace() const
Module * getParent()
Get the module that this global value is contained inside of...
@ InternalLinkage
Rename collisions when linking (static functions).
static LLVM_ABI IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
constexpr bool isValid() const
constexpr uint16_t getNumElements() const
Returns the number of elements in a vector LLT.
constexpr bool isVector() const
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
int getNumber() const
MachineBasicBlocks are uniquely numbered at the function level, unless they're not in a MachineFuncti...
LLVM_ABI iterator getFirstNonPHI()
Returns a pointer to the first instruction in this block that is not a PHINode instruction.
const MachineFunction * getParent() const
Return the MachineFunction containing this basic block.
MachineInstrBundleIterator< MachineInstr > iterator
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
Helper class to build MachineInstr.
MachineBasicBlock::iterator getInsertPt()
Current insertion point for new instructions.
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
MachineFunction & getMF()
Getter for the function we currently build.
MachineRegisterInfo * getMRI()
Getter for MRI.
void constrainAllUses(const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI) const
const MachineInstrBuilder & addUse(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register use operand.
const MachineInstrBuilder & addReg(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a new virtual register operand.
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
const MachineInstrBuilder & add(const MachineOperand &MO) const
const MachineInstrBuilder & addMBB(MachineBasicBlock *MBB, unsigned TargetFlags=0) const
const MachineInstrBuilder & addDef(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register definition operand.
const MachineInstrBuilder & setMIFlags(unsigned Flags) const
MachineInstr * getInstr() const
If conversion operators fail, use this method to get the MachineInstr explicitly.
Representation of each machine instruction.
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
const MachineBasicBlock * getParent() const
unsigned getNumOperands() const
Retuns the total number of operands.
LLVM_ABI unsigned getNumExplicitOperands() const
Returns the number of non-implicit operands.
LLVM_ABI unsigned getNumExplicitDefs() const
Returns the number of non-implicit definitions.
LLVM_ABI void emitGenericError(const Twine &ErrMsg) const
LLVM_ABI const MachineFunction * getMF() const
Return the function that contains the basic block that this instruction belongs to.
const DebugLoc & getDebugLoc() const
Returns the debug location id of this MachineInstr.
const MachineOperand & getOperand(unsigned i) const
A description of a memory reference used in the backend.
@ MOVolatile
The memory access is volatile.
@ MONonTemporal
The memory access is non-temporal.
bool isReg() const
isReg - Tests if this is a MO_Register operand.
MachineBasicBlock * getMBB() const
Register getReg() const
getReg - Returns the register number.
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
defusechain_instr_iterator< true, false, false, true > use_instr_iterator
use_instr_iterator/use_instr_begin/use_instr_end - Walk all uses of the specified register,...
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
LLVM_ABI MachineInstr * getVRegDef(Register Reg) const
getVRegDef - Return the machine instr that defines the specified virtual register or null if none is ...
use_instr_iterator use_instr_begin(Register RegNo) const
bool use_nodbg_empty(Register RegNo) const
use_nodbg_empty - Return true if there are no non-Debug instructions using the specified register.
static def_instr_iterator def_instr_end()
defusechain_instr_iterator< false, true, false, true > def_instr_iterator
def_instr_iterator/def_instr_begin/def_instr_end - Walk all defs of the specified register,...
LLVM_ABI Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
def_instr_iterator def_instr_begin(Register RegNo) const
LLT getType(Register Reg) const
Get the low-level type of Reg or LLT{} if Reg is not a generic (target independent) virtual register.
static use_instr_iterator use_instr_end()
iterator_range< use_instr_nodbg_iterator > use_nodbg_instructions(Register Reg) const
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
const MachineFunction & getMF() const
LLVM_ABI void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
LLVM_ABI Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
const TargetRegisterClass * getRegClassOrNull(Register Reg) const
Return the register class of Reg, or null if Reg has not been assigned a register class yet.
iterator_range< use_instr_iterator > use_instructions(Register Reg) const
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
LLVM_ABI void replaceRegWith(Register FromReg, Register ToReg)
replaceRegWith - Replace all instances of FromReg with ToReg in the machine function.
Analysis providing profile information.
Holds all the information related to register banks.
Wrapper class representing virtual and physical registers.
constexpr bool isValid() const
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
bool isScalarOrVectorSigned(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateOpTypeSampledImage(SPIRVTypeInst ImageType, MachineIRBuilder &MIRBuilder)
void assignSPIRVTypeToVReg(SPIRVTypeInst Type, Register VReg, const MachineFunction &MF)
const TargetRegisterClass * getRegClass(SPIRVTypeInst SpvType) const
MachineInstr * getOrAddMemAliasingINTELInst(MachineIRBuilder &MIRBuilder, const MDNode *AliasingListMD)
bool isAggregateType(SPIRVTypeInst Type) const
unsigned getScalarOrVectorBitWidth(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst getOrCreateSPIRVVectorType(SPIRVTypeInst BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
Register buildGlobalVariable(Register Reg, SPIRVTypeInst BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
SPIRVTypeInst getResultType(Register VReg, MachineFunction *MF=nullptr)
unsigned getScalarOrVectorComponentCount(Register VReg) const
const Type * getTypeForSPIRVType(SPIRVTypeInst Ty) const
bool isBitcastCompatible(SPIRVTypeInst Type1, SPIRVTypeInst Type2) const
unsigned getPointerSize() const
Register getOrCreateConstFP(APFloat Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
LLT getRegType(SPIRVTypeInst SpvType) const
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC)
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
Register getSPIRVTypeID(SPIRVTypeInst SpirvType) const
Register getOrCreateConstInt(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
bool findValueAttrs(const MachineInstr *Key, Type *&Ty, StringRef &Name)
SPIRVTypeInst retrieveScalarOrVectorIntType(SPIRVTypeInst Type) const
Register getOrCreateGlobalVariableWithBinding(SPIRVTypeInst VarType, uint32_t Set, uint32_t Binding, StringRef Name, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst changePointerStorageClass(SPIRVTypeInst PtrType, SPIRV::StorageClass::StorageClass SC, MachineInstr &I)
Register getOrCreateConstVector(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType=nullptr)
void addGlobalObject(const Value *V, const MachineFunction *MF, Register R)
SPIRVTypeInst getScalarOrVectorComponentType(SPIRVTypeInst Type) const
void recordFunctionPointer(const MachineOperand *MO, const Function *F)
SPIRVTypeInst getOrCreateSPIRVFloatType(unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII)
SPIRVTypeInst getPointeeType(SPIRVTypeInst PtrType)
SPIRVTypeInst getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
MachineFunction * setCurrentFunc(MachineFunction &MF)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType)
SPIRVTypeInst getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
Type * getDeducedGlobalValueType(const GlobalValue *Global)
Register getOrCreateUndef(MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
bool erase(const MachineInstr *MI)
bool add(SPIRV::IRHandle Handle, const MachineInstr *MI)
Register find(SPIRV::IRHandle Handle, const MachineFunction *MF)
bool isPhysicalSPIRV() const
bool isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const
bool canUseExtInstSet(SPIRV::InstructionSet::InstructionSet E) const
bool isLogicalSPIRV() const
bool canUseExtension(SPIRV::Extension::Extension E) const
bool isTypeIntOrFloat() 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
Get the string size.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
@ HalfTyID
16-bit floating point type
@ FloatTyID
32-bit floating point type
@ BFloatTyID
16-bit floating point type (7-bit significand)
@ 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.
self_iterator getIterator()
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.
constexpr std::underlying_type_t< E > Mask()
Get a bitmask with 1s in all places up to the high-order bit of E's largest value.
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
NodeAddr< FuncNode * > Func
BaseReg
Stack frame base register. Bit 0 of FREInfo.Info.
This is an optimization pass for GlobalISel generic memory operations.
@ Low
Lower the current thread's priority such that it does not affect foreground tasks significantly.
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
void addStringImm(StringRef Str, MCInst &Inst)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineFunction &MF)
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.
MachineInstr * getDef(const MachineOperand &MO, const MachineRegisterInfo *MRI)
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)
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, ArrayRef< uint32_t > DecArgs, StringRef StrImm)
LLVM_ABI bool isNullOrNullSplat(const MachineInstr &MI, const MachineRegisterInfo &MRI, bool AllowUndefs=false)
Return true if the value is a constant 0 integer or a splatted vector of a constant 0 integer (with n...
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)
RelativeUniformCounterPtr ValuesPtrExpr VTableAddr Value
void buildOpName(Register Target, StringRef Name, MachineIRBuilder &MIRBuilder)
bool any_of(R &&range, UnaryPredicate P)
Provide wrappers to std::any_of which take ranges instead of having to pass begin/end explicitly.
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
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)
bool isSpvIntrinsic(const MachineInstr &MI, Intrinsic::ID IntrinsicID)
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...
MCRegisterClass TargetRegisterClass