37#include "llvm/IR/IntrinsicsSPIRV.h"
43#define DEBUG_TYPE "spirv-isel"
50 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
55 std::optional<Register> Bias;
56 std::optional<Register>
Offset;
57 std::optional<Register> MinLod;
58 std::optional<Register> GradX;
59 std::optional<Register> GradY;
60 std::optional<Register> Lod;
61 std::optional<Register> Compare;
64llvm::SPIRV::SelectionControl::SelectionControl
65getSelectionOperandForImm(
int Imm) {
67 return SPIRV::SelectionControl::Flatten;
69 return SPIRV::SelectionControl::DontFlatten;
71 return SPIRV::SelectionControl::None;
75#define GET_GLOBALISEL_PREDICATE_BITSET
76#include "SPIRVGenGlobalISel.inc"
77#undef GET_GLOBALISEL_PREDICATE_BITSET
104#define GET_GLOBALISEL_PREDICATES_DECL
105#include "SPIRVGenGlobalISel.inc"
106#undef GET_GLOBALISEL_PREDICATES_DECL
108#define GET_GLOBALISEL_TEMPORARIES_DECL
109#include "SPIRVGenGlobalISel.inc"
110#undef GET_GLOBALISEL_TEMPORARIES_DECL
134 unsigned BitSetOpcode)
const;
138 unsigned BitSetOpcode)
const;
142 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
149 unsigned Opcode)
const;
152 unsigned Opcode)
const;
174 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
183 Register castPtrToPtrToInt(
Register Ptr, SPIRV::StorageClass::StorageClass SC,
187 bool selectAtomicPtrValue(
203 unsigned OpType)
const;
271 unsigned Opcode)
const;
275 unsigned Opcode)
const;
279 unsigned Opcode)
const;
283 unsigned Opcode)
const;
285 template <
bool Signed>
288 template <
bool Signed>
295 template <
typename PickOpcodeFn>
298 PickOpcodeFn &&PickOpcode)
const;
315 template <
typename PickOpcodeFn>
318 PickOpcodeFn &&PickOpcode)
const;
336 bool IsSigned)
const;
338 bool IsSigned,
unsigned Opcode)
const;
340 bool IsSigned)
const;
346 bool IsSigned)
const;
387 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
388 bool useMISrc =
true,
390 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
391 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
392 bool useMISrc =
true,
394 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
395 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
396 bool setMIFlags =
true,
bool useMISrc =
true,
398 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
399 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
400 bool useMISrc =
true,
403 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
404 MachineInstr &
I)
const;
406 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
407 MachineInstr &
I)
const;
409 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
412 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I,
unsigned Opcode)
const;
415 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
416 bool WithGroupSync)
const;
418 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
419 MachineInstr &
I)
const;
421 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
422 MachineInstr &
I)
const;
426 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
427 MachineInstr &
I)
const;
429 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
430 MachineInstr &
I)
const;
432 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
433 MachineInstr &
I)
const;
434 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
435 MachineInstr &
I)
const;
436 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
437 SPIRVTypeInst ResType,
438 MachineInstr &
I)
const;
439 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
440 MachineInstr &
I)
const;
443 std::optional<Register> LodReg = std::nullopt)
const;
444 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
445 MachineInstr &
I)
const;
446 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
447 MachineInstr &
I)
const;
448 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
449 MachineInstr &
I)
const;
450 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
451 MachineInstr &
I)
const;
452 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
453 MachineInstr &
I)
const;
454 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
455 MachineInstr &
I)
const;
456 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
457 MachineInstr &
I)
const;
458 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
459 SPIRVTypeInst ResType,
460 MachineInstr &
I)
const;
461 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
462 MachineInstr &
I)
const;
463 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
464 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
465 MachineInstr &
I)
const;
466 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
467 MachineInstr &
I)
const;
468 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
469 MachineInstr &
I)
const;
470 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
471 MachineInstr &
I)
const;
472 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
473 MachineInstr &
I)
const;
474 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
475 MachineInstr &
I)
const;
477 bool selectCopySign(
Register ResVReg, SPIRVTypeInst ResType,
478 MachineInstr &
I)
const;
480 bool selectLdexp(
Register ResVReg, SPIRVTypeInst ResType,
481 MachineInstr &
I)
const;
482 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
483 MachineInstr &
I)
const;
484 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
485 MachineInstr &
I)
const;
486 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
487 MachineInstr &
I,
const unsigned DPdOpCode)
const;
489 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
490 SPIRVTypeInst ResType =
nullptr)
const;
491 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
492 SPIRVTypeInst ResType =
nullptr)
const;
494 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
495 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
496 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
498 MachineInstr &
I)
const;
499 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
501 bool wrapIntoSpecConstantOp(MachineInstr &
I,
504 Register getUcharPtrTypeReg(MachineInstr &
I,
505 SPIRV::StorageClass::StorageClass SC)
const;
506 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
508 uint32_t Opcode)
const;
509 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
510 SPIRVTypeInst SrcPtrTy)
const;
511 Register buildPointerToResource(SPIRVTypeInst ResType,
512 SPIRV::StorageClass::StorageClass SC,
513 uint32_t Set, uint32_t
Binding,
514 uint32_t ArraySize,
Register IndexReg,
516 MachineIRBuilder MIRBuilder)
const;
517 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
518 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
519 Register &ReadReg, MachineInstr &InsertionPoint)
const;
520 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
523 const ImageOperands *ImOps =
nullptr)
const;
524 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
526 Register CoordinateReg,
const ImageOperands &ImOps,
529 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
530 Register ResVReg, SPIRVTypeInst ResType,
531 MachineInstr &
I)
const;
532 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
533 Register ResVReg, SPIRVTypeInst ResType,
534 MachineInstr &
I)
const;
535 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
536 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
537 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
538 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
541 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
542 Register SrcReg,
unsigned int Opcode,
543 std::function<
bool(
Register, SPIRVTypeInst,
544 MachineInstr &,
Register,
unsigned)>
548bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
550 if (
TET->getTargetExtName() ==
"spirv.Image") {
553 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
554 return TET->getTypeParameter(0)->isIntegerTy();
558#define GET_GLOBALISEL_IMPL
559#include "SPIRVGenGlobalISel.inc"
560#undef GET_GLOBALISEL_IMPL
566 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
569#include
"SPIRVGenGlobalISel.inc"
572#include
"SPIRVGenGlobalISel.inc"
584 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
589 if (HasVRegsReset == &MF)
604 for (
const auto &
MBB : MF) {
605 for (
const auto &
MI :
MBB) {
608 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
612 LLT DstType = MRI.
getType(DstReg);
614 LLT SrcType = MRI.
getType(SrcReg);
615 if (DstType != SrcType)
620 if (DstRC != SrcRC && SrcRC)
632 while (!Stack.empty()) {
637 switch (
MI->getOpcode()) {
638 case TargetOpcode::G_INTRINSIC:
639 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
640 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
643 if (IntrID != Intrinsic::spv_const_composite &&
644 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
648 case TargetOpcode::G_BUILD_VECTOR:
649 case TargetOpcode::G_SPLAT_VECTOR:
651 i < OpDef->getNumOperands(); i++) {
656 Stack.push_back(OpNestedDef);
659 case TargetOpcode::G_CONSTANT:
660 case TargetOpcode::G_FCONSTANT:
661 case TargetOpcode::G_IMPLICIT_DEF:
662 case SPIRV::OpConstantTrue:
663 case SPIRV::OpConstantFalse:
664 case SPIRV::OpConstantI:
665 case SPIRV::OpConstantF:
666 case SPIRV::OpConstantComposite:
667 case SPIRV::OpConstantCompositeContinuedINTEL:
668 case SPIRV::OpConstantSampler:
669 case SPIRV::OpConstantNull:
671 case SPIRV::OpPoisonKHR:
672 case SPIRV::OpConstantFunctionPointerINTEL:
699 case Intrinsic::spv_all:
700 case Intrinsic::spv_alloca:
701 case Intrinsic::spv_any:
702 case Intrinsic::spv_bitcast:
703 case Intrinsic::spv_const_composite:
704 case Intrinsic::spv_degrees:
705 case Intrinsic::spv_distance:
706 case Intrinsic::spv_extractelt:
707 case Intrinsic::spv_extractv:
708 case Intrinsic::spv_faceforward:
709 case Intrinsic::spv_fdot:
710 case Intrinsic::spv_firstbitlow:
711 case Intrinsic::spv_firstbitshigh:
712 case Intrinsic::spv_firstbituhigh:
713 case Intrinsic::spv_frac:
714 case Intrinsic::spv_gep:
715 case Intrinsic::spv_global_offset:
716 case Intrinsic::spv_global_size:
717 case Intrinsic::spv_group_id:
718 case Intrinsic::spv_insertelt:
719 case Intrinsic::spv_insertv:
720 case Intrinsic::spv_isinf:
721 case Intrinsic::spv_isnan:
722 case Intrinsic::spv_isfinite:
723 case Intrinsic::spv_isnormal:
724 case Intrinsic::spv_lerp:
725 case Intrinsic::spv_length:
726 case Intrinsic::spv_normalize:
727 case Intrinsic::spv_num_subgroups:
728 case Intrinsic::spv_num_workgroups:
729 case Intrinsic::spv_ptrcast:
730 case Intrinsic::spv_radians:
731 case Intrinsic::spv_reflect:
732 case Intrinsic::spv_refract:
733 case Intrinsic::spv_resource_getbasepointer:
734 case Intrinsic::spv_resource_getpointer:
735 case Intrinsic::spv_resource_handlefrombinding:
736 case Intrinsic::spv_resource_handlefromimplicitbinding:
737 case Intrinsic::spv_resource_nonuniformindex:
738 case Intrinsic::spv_resource_sample:
739 case Intrinsic::spv_rsqrt:
740 case Intrinsic::spv_saturate:
741 case Intrinsic::spv_sdot:
742 case Intrinsic::spv_sign:
743 case Intrinsic::spv_smoothstep:
744 case Intrinsic::spv_subgroup_id:
745 case Intrinsic::spv_subgroup_local_invocation_id:
746 case Intrinsic::spv_subgroup_max_size:
747 case Intrinsic::spv_subgroup_size:
748 case Intrinsic::spv_thread_id:
749 case Intrinsic::spv_thread_id_in_group:
750 case Intrinsic::spv_udot:
751 case Intrinsic::spv_undef:
752 case Intrinsic::spv_value_md:
753 case Intrinsic::spv_workgroup_size:
765 case SPIRV::OpTypeVoid:
766 case SPIRV::OpTypeBool:
767 case SPIRV::OpTypeInt:
768 case SPIRV::OpTypeFloat:
769 case SPIRV::OpTypeVector:
770 case SPIRV::OpTypeVectorIdEXT:
771 case SPIRV::OpTypeMatrix:
772 case SPIRV::OpTypeImage:
773 case SPIRV::OpTypeSampler:
774 case SPIRV::OpTypeSampledImage:
775 case SPIRV::OpTypeArray:
776 case SPIRV::OpTypeRuntimeArray:
777 case SPIRV::OpTypeStruct:
778 case SPIRV::OpTypeOpaque:
779 case SPIRV::OpTypePointer:
780 case SPIRV::OpTypeFunction:
781 case SPIRV::OpTypeEvent:
782 case SPIRV::OpTypeDeviceEvent:
783 case SPIRV::OpTypeReserveId:
784 case SPIRV::OpTypeQueue:
785 case SPIRV::OpTypePipe:
786 case SPIRV::OpTypeForwardPointer:
787 case SPIRV::OpTypePipeStorage:
788 case SPIRV::OpTypeNamedBarrier:
789 case SPIRV::OpTypeAccelerationStructureNV:
790 case SPIRV::OpTypeCooperativeMatrixNV:
791 case SPIRV::OpTypeCooperativeMatrixKHR:
801 if (
MI.getNumDefs() == 0)
804 for (
const auto &MO :
MI.all_defs()) {
806 if (
Reg.isPhysical()) {
811 if (
UseMI.getOpcode() != SPIRV::OpName) {
818 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
819 MI.isLifetimeMarker()) {
822 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
833 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
834 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
837 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
842 if (
MI.mayStore() ||
MI.isCall() ||
843 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
844 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
845 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
856 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
863void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
865 for (
const auto &MO :
MI.all_defs()) {
869 SmallVector<MachineInstr *, 4> UselessOpNames;
872 "There is still a use of the dead function.");
875 for (MachineInstr *OpNameMI : UselessOpNames) {
877 OpNameMI->eraseFromParent();
882void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
885 removeOpNamesForDeadMI(
MI);
886 MI.eraseFromParent();
889bool SPIRVInstructionSelector::select(MachineInstr &
I) {
890 resetVRegsType(*
I.getParent()->getParent());
892 assert(
I.getParent() &&
"Instruction should be in a basic block!");
893 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
898 removeDeadInstruction(
I);
905 if (Opcode == SPIRV::ASSIGN_TYPE) {
906 Register DstReg =
I.getOperand(0).getReg();
907 Register SrcReg =
I.getOperand(1).getReg();
910 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
911 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
912 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
913 Register SelectDstReg =
Def->getOperand(0).getReg();
914 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
916 assert(SuccessToSelectSelect);
918 Def->eraseFromParent();
925 bool Res = selectImpl(
I, *CoverageInfo);
927 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
928 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
932 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
944 }
else if (
I.getNumDefs() == 1) {
956 removeDeadInstruction(
I);
961 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
962 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
968 bool HasDefs =
I.getNumDefs() > 0;
971 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
972 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
973 if (spvSelect(ResVReg, ResType,
I)) {
975 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
986 case TargetOpcode::G_CONSTANT:
987 case TargetOpcode::G_FCONSTANT:
994 MachineInstr &
I)
const {
997 if (DstRC != SrcRC && SrcRC)
999 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
1006bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
1007 SPIRVTypeInst ResType,
1008 MachineInstr &
I)
const {
1009 const unsigned Opcode =
I.getOpcode();
1011 return selectImpl(
I, *CoverageInfo);
1013 case TargetOpcode::G_CONSTANT:
1014 case TargetOpcode::G_FCONSTANT:
1015 return selectConst(ResVReg, ResType,
I);
1016 case TargetOpcode::G_GLOBAL_VALUE:
1017 return selectGlobalValue(ResVReg,
I);
1018 case TargetOpcode::G_IMPLICIT_DEF:
1019 return selectOpUndef(ResVReg, ResType,
I);
1020 case TargetOpcode::G_FREEZE:
1021 return selectFreeze(ResVReg, ResType,
I);
1023 case TargetOpcode::G_INTRINSIC:
1024 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1025 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1026 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1027 return selectIntrinsic(ResVReg, ResType,
I);
1028 case TargetOpcode::G_BITREVERSE:
1029 return selectBitreverse(ResVReg, ResType,
I);
1031 case TargetOpcode::G_BUILD_VECTOR:
1032 return selectBuildVector(ResVReg, ResType,
I);
1033 case TargetOpcode::G_SPLAT_VECTOR:
1034 return selectSplatVector(ResVReg, ResType,
I);
1035 case TargetOpcode::G_CONCAT_VECTORS:
1036 return selectConcatVectors(ResVReg, ResType,
I);
1038 case TargetOpcode::G_SHUFFLE_VECTOR: {
1039 MachineBasicBlock &BB = *
I.getParent();
1040 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1043 .
addUse(
I.getOperand(1).getReg())
1044 .
addUse(
I.getOperand(2).getReg());
1045 for (
auto V :
I.getOperand(3).getShuffleMask())
1050 case TargetOpcode::G_MEMMOVE:
1051 case TargetOpcode::G_MEMCPY:
1052 case TargetOpcode::G_MEMCPY_INLINE:
1053 case TargetOpcode::G_MEMSET:
1054 case TargetOpcode::G_MEMSET_INLINE:
1055 return selectMemOperation(ResVReg,
I);
1057 case TargetOpcode::G_ICMP:
1058 return selectICmp(ResVReg, ResType,
I);
1059 case TargetOpcode::G_FCMP:
1060 return selectFCmp(ResVReg, ResType,
I);
1062 case TargetOpcode::G_FRAME_INDEX:
1063 return selectFrameIndex(ResVReg, ResType,
I);
1065 case TargetOpcode::G_LOAD:
1066 return selectLoad(ResVReg, ResType,
I);
1067 case TargetOpcode::G_STORE:
1068 return selectStore(
I);
1070 case TargetOpcode::G_BR:
1071 return selectBranch(
I);
1072 case TargetOpcode::G_BRCOND:
1073 return selectBranchCond(
I);
1075 case TargetOpcode::G_PHI:
1076 return selectPhi(ResVReg,
I);
1078 case TargetOpcode::G_FPTOSI:
1079 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1080 case TargetOpcode::G_FPTOUI:
1081 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1083 case TargetOpcode::G_FPTOSI_SAT:
1084 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1085 case TargetOpcode::G_FPTOUI_SAT:
1086 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1088 case TargetOpcode::G_SITOFP:
1089 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1090 case TargetOpcode::G_UITOFP:
1091 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1093 case TargetOpcode::G_CTPOP:
1094 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1095 case TargetOpcode::G_SMIN:
1096 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1097 case TargetOpcode::G_UMIN:
1098 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1100 case TargetOpcode::G_SMAX:
1101 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1102 case TargetOpcode::G_UMAX:
1103 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1105 case TargetOpcode::G_SCMP:
1106 return selectSUCmp(ResVReg, ResType,
I,
true);
1107 case TargetOpcode::G_UCMP:
1108 return selectSUCmp(ResVReg, ResType,
I,
false);
1109 case TargetOpcode::G_LROUND:
1110 case TargetOpcode::G_LLROUND: {
1113 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1115 regForLround, *(
I.getParent()->getParent()));
1117 CL::round, GL::Round,
false);
1119 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1126 case TargetOpcode::G_STRICT_FMA:
1127 case TargetOpcode::G_FMA: {
1130 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1133 .
addUse(
I.getOperand(1).getReg())
1134 .
addUse(
I.getOperand(2).getReg())
1135 .
addUse(
I.getOperand(3).getReg())
1140 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1143 case TargetOpcode::G_FLDEXP:
1144 case TargetOpcode::G_STRICT_FLDEXP:
1145 return selectLdexp(ResVReg, ResType,
I);
1147 case TargetOpcode::G_FPOW:
1148 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1149 case TargetOpcode::G_FPOWI:
1150 return selectFpowi(ResVReg, ResType,
I);
1152 case TargetOpcode::G_FEXP:
1153 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1154 case TargetOpcode::G_FEXP2:
1155 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1156 case TargetOpcode::G_FEXP10:
1157 return selectExp10(ResVReg, ResType,
I);
1159 case TargetOpcode::G_FMODF:
1160 return selectModf(ResVReg, ResType,
I);
1161 case TargetOpcode::G_FSINCOS:
1162 return selectSincos(ResVReg, ResType,
I);
1164 case TargetOpcode::G_FLOG:
1165 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1166 case TargetOpcode::G_FLOG2:
1167 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1168 case TargetOpcode::G_FLOG10:
1169 return selectLog10(ResVReg, ResType,
I);
1171 case TargetOpcode::G_FABS:
1172 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1173 case TargetOpcode::G_ABS:
1174 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1176 case TargetOpcode::G_FMINNUM:
1177 case TargetOpcode::G_FMINIMUM:
1178 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1179 case TargetOpcode::G_FMAXNUM:
1180 case TargetOpcode::G_FMAXIMUM:
1181 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1183 case TargetOpcode::G_FCOPYSIGN:
1184 return selectCopySign(ResVReg, ResType,
I);
1186 case TargetOpcode::G_FCEIL:
1187 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1188 case TargetOpcode::G_FFLOOR:
1189 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1191 case TargetOpcode::G_FCOS:
1192 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1193 case TargetOpcode::G_FSIN:
1194 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1195 case TargetOpcode::G_FTAN:
1196 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1197 case TargetOpcode::G_FACOS:
1198 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1199 case TargetOpcode::G_FASIN:
1200 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1201 case TargetOpcode::G_FATAN:
1202 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1203 case TargetOpcode::G_FATAN2:
1204 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1205 case TargetOpcode::G_FCOSH:
1206 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1207 case TargetOpcode::G_FSINH:
1208 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1209 case TargetOpcode::G_FTANH:
1210 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1212 case TargetOpcode::G_STRICT_FSQRT:
1213 case TargetOpcode::G_FSQRT:
1214 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1216 case TargetOpcode::G_CTTZ:
1217 case TargetOpcode::G_CTTZ_ZERO_POISON:
1218 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1219 case TargetOpcode::G_CTLZ:
1220 case TargetOpcode::G_CTLZ_ZERO_POISON:
1221 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1223 case TargetOpcode::G_INTRINSIC_ROUND:
1224 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1225 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1226 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1227 case TargetOpcode::G_INTRINSIC_TRUNC:
1228 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1229 case TargetOpcode::G_FRINT:
1230 case TargetOpcode::G_FNEARBYINT:
1231 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1233 case TargetOpcode::G_SMULH:
1234 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1235 case TargetOpcode::G_UMULH:
1236 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1238 case TargetOpcode::G_SADDSAT:
1239 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1240 case TargetOpcode::G_UADDSAT:
1241 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1242 case TargetOpcode::G_SSUBSAT:
1243 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1244 case TargetOpcode::G_USUBSAT:
1245 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1247 case TargetOpcode::G_FFREXP:
1248 return selectFrexp(ResVReg, ResType,
I);
1250 case TargetOpcode::G_UADDO:
1251 return selectOverflowArith(ResVReg, ResType,
I,
1253 : SPIRV::OpIAddCarryS);
1254 case TargetOpcode::G_USUBO:
1255 return selectOverflowArith(ResVReg, ResType,
I,
1257 : SPIRV::OpISubBorrowS);
1258 case TargetOpcode::G_UMULO:
1259 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1260 case TargetOpcode::G_SMULO:
1261 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1263 case TargetOpcode::G_SEXT:
1264 return selectExt(ResVReg, ResType,
I,
true);
1265 case TargetOpcode::G_ANYEXT:
1266 case TargetOpcode::G_ZEXT:
1267 return selectExt(ResVReg, ResType,
I,
false);
1268 case TargetOpcode::G_TRUNC:
1269 return selectTrunc(ResVReg, ResType,
I);
1270 case TargetOpcode::G_FPTRUNC:
1271 case TargetOpcode::G_FPEXT:
1272 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1274 case TargetOpcode::G_PTRTOINT:
1275 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1276 case TargetOpcode::G_INTTOPTR:
1277 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1278 case TargetOpcode::G_BITCAST:
1279 return selectBitcast(ResVReg, ResType,
I);
1280 case TargetOpcode::G_ADDRSPACE_CAST:
1281 return selectAddrSpaceCast(ResVReg, ResType,
I);
1282 case TargetOpcode::G_PTRMASK:
1283 return selectPtrMask(ResVReg, ResType,
I);
1284 case TargetOpcode::G_PTR_ADD: {
1286 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1290 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1291 (*II).getOpcode() == TargetOpcode::COPY ||
1292 (*II).getOpcode() == SPIRV::OpVariable ||
1293 (*II).getOpcode() == SPIRV::OpUntypedVariableKHR) &&
1294 getImm(
I.getOperand(2), MRI));
1296 bool IsGVInit =
false;
1300 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1301 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1302 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1303 (*UseIt).getOpcode() == SPIRV::OpVariable ||
1304 (*UseIt).getOpcode() == SPIRV::OpUntypedVariableKHR) {
1316 const bool UseUntypedPointers =
1317 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1318 if (UseUntypedPointers) {
1319 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1322 .
addImm(
static_cast<uint32_t
>(
1323 SPIRV::Opcode::UntypedInBoundsPtrAccessChainKHR))
1326 .
addUse(
I.getOperand(2).getReg())
1333 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1345 return diagnoseUnsupported(
1346 I,
"incompatible result and operand types in a bitcast");
1348 MachineInstrBuilder MIB =
1349 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1356 : SPIRV::OpInBoundsPtrAccessChain))
1360 .
addUse(
I.getOperand(2).getReg())
1363 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1367 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1369 .
addUse(
I.getOperand(2).getReg())
1378 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1381 .
addImm(
static_cast<uint32_t
>(
1382 SPIRV::Opcode::InBoundsPtrAccessChain))
1385 .
addUse(
I.getOperand(2).getReg());
1390 case TargetOpcode::G_ATOMICRMW_OR:
1391 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1392 case TargetOpcode::G_ATOMICRMW_ADD:
1393 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1394 case TargetOpcode::G_ATOMICRMW_AND:
1395 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1396 case TargetOpcode::G_ATOMICRMW_MAX:
1397 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1398 case TargetOpcode::G_ATOMICRMW_MIN:
1399 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1400 case TargetOpcode::G_ATOMICRMW_SUB:
1401 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1402 case TargetOpcode::G_ATOMICRMW_XOR:
1403 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1404 case TargetOpcode::G_ATOMICRMW_UMAX:
1405 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1406 case TargetOpcode::G_ATOMICRMW_UMIN:
1407 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1408 case TargetOpcode::G_ATOMICRMW_XCHG:
1409 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1411 case TargetOpcode::G_ATOMICRMW_FADD:
1412 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1413 case TargetOpcode::G_ATOMICRMW_FSUB:
1415 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1417 : SPIRV::OpFNegate);
1418 case TargetOpcode::G_ATOMICRMW_FMIN:
1419 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1420 case TargetOpcode::G_ATOMICRMW_FMAX:
1421 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1423 case TargetOpcode::G_FENCE:
1424 return selectFence(
I);
1426 case TargetOpcode::G_STACKSAVE:
1427 return selectStackSave(ResVReg, ResType,
I);
1428 case TargetOpcode::G_STACKRESTORE:
1429 return selectStackRestore(
I);
1431 case TargetOpcode::G_UNMERGE_VALUES:
1434 case TargetOpcode::G_TRAP:
1435 case TargetOpcode::G_UBSANTRAP:
1436 return selectTrap(
I);
1441 case TargetOpcode::DBG_LABEL:
1443 case TargetOpcode::G_DEBUGTRAP:
1444 return selectDebugTrap(ResVReg, ResType,
I);
1445 case TargetOpcode::G_PREFETCH:
1446 return selectPrefetch(
I);
1453bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1454 SPIRVTypeInst ResType,
1455 MachineInstr &
I)
const {
1456 unsigned Opcode = SPIRV::OpNop;
1463bool SPIRVInstructionSelector::selectPrefetch(MachineInstr &
I)
const {
1472 MachineIRBuilder MIRBuilder(
I);
1474 const SPIRVTypeInst PointerSizeType =
1482 Register AddrVal =
I.getOperand(0).getReg();
1485 return selectExtInst(ExtReg, GR.
getOpTypeVoid(MIRBuilder),
I, CL::prefetch,
1487 {AddrVal, ConstIntOne});
1492bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1493 SPIRVTypeInst ResType,
1495 GL::GLSLExtInst GLInst,
1496 bool setMIFlags,
bool useMISrc,
1499 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1500 return diagnoseUnsupported(
1502 "this instruction is only supported with the GLSL extended instruction "
1504 return selectExtInst(ResVReg, ResType,
I,
1505 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1506 setMIFlags, useMISrc, SrcRegs);
1509bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1510 SPIRVTypeInst ResType,
1512 CL::OpenCLExtInst CLInst,
1513 bool setMIFlags,
bool useMISrc,
1515 return selectExtInst(ResVReg, ResType,
I,
1516 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1517 setMIFlags, useMISrc, SrcRegs);
1520bool SPIRVInstructionSelector::selectExtInst(
1521 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1522 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1524 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1525 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1526 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1530bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1531 SPIRVTypeInst ResType,
1534 bool setMIFlags,
bool useMISrc,
1537 for (
const auto &[InstructionSet, Opcode] : Insts) {
1541 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1544 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1549 const unsigned NumOps =
I.getNumOperands();
1552 I.getOperand(Index).getType() ==
1553 MachineOperand::MachineOperandType::MO_IntrinsicID)
1556 MIB.
add(
I.getOperand(Index));
1568bool SPIRVInstructionSelector::selectCopySign(
Register ResVReg,
1569 SPIRVTypeInst ResType,
1570 MachineInstr &
I)
const {
1572 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1577 Register MagnitudeReg =
I.getOperand(1).getReg();
1578 Register SignReg =
I.getOperand(2).getReg();
1586 unsigned AndOpcode, OrOpcode;
1587 if (ComponentCount > 1) {
1591 AndOpcode = SPIRV::OpBitwiseAndV;
1592 OrOpcode = SPIRV::OpBitwiseOrV;
1596 AndOpcode = SPIRV::OpBitwiseAndS;
1597 OrOpcode = SPIRV::OpBitwiseOrS;
1603 return selectOpWithSrcs(ResReg, IntType,
I, SrcRegs, Opcode);
1606 Register MagnitudeInt, SignInt, MagnitudeBits, SignBits, CombinedInt;
1607 if (!EmitBitOp(MagnitudeInt, {MagnitudeReg}, SPIRV::OpBitcast) ||
1608 !EmitBitOp(SignInt, {SignReg}, SPIRV::OpBitcast) ||
1609 !EmitBitOp(MagnitudeBits, {MagnitudeInt, NotSignMask}, AndOpcode) ||
1610 !EmitBitOp(SignBits, {SignInt, SignMask}, AndOpcode) ||
1611 !EmitBitOp(CombinedInt, {MagnitudeBits, SignBits}, OrOpcode))
1614 return selectOpWithSrcs(ResVReg, ResType,
I, {CombinedInt}, SPIRV::OpBitcast);
1617bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1618 SPIRVTypeInst ResType,
1619 MachineInstr &
I)
const {
1620 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1621 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1622 for (
const auto &Ex : ExtInsts) {
1623 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1624 uint32_t Opcode = Ex.second;
1628 MachineIRBuilder MIRBuilder(
I);
1631 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1638 const bool IsUntyped =
1639 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1641 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1642 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1643 : SPIRV::OpVariable))
1646 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1652 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1655 .
addImm(
static_cast<uint32_t
>(Ex.first))
1657 .
add(
I.getOperand(2))
1661 Register ExpResReg =
I.getOperand(1).getReg();
1663 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1673bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1674 SPIRVTypeInst ResType,
1675 MachineInstr &
I)
const {
1676 Register XReg =
I.getOperand(1).getReg();
1677 Register ExpReg =
I.getOperand(2).getReg();
1683 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1684 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1686 SPIRVTypeInst ExpVecType =
1690 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1691 TII.get(SPIRV::OpCompositeConstruct))
1694 for (
unsigned J = 0; J < NumElts; ++J)
1700 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1701 true,
false, {XReg, ExpReg});
1704bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1705 SPIRVTypeInst ResType,
1706 MachineInstr &
I)
const {
1707 Register CosResVReg =
I.getOperand(1).getReg();
1708 unsigned SrcIdx =
I.getNumExplicitDefs();
1713 MachineIRBuilder MIRBuilder(
I);
1715 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1722 const bool IsUntyped =
1723 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1725 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1726 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1727 : SPIRV::OpVariable))
1730 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1734 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1737 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1739 .
add(
I.getOperand(SrcIdx))
1743 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1751 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1754 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1756 .
add(
I.getOperand(SrcIdx))
1758 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1761 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1763 .
add(
I.getOperand(SrcIdx))
1770bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1771 SPIRVTypeInst ResType,
1774 unsigned Opcode)
const {
1775 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1785bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1786 SPIRVTypeInst ResType,
1789 unsigned Opcode)
const {
1790 MachineIRBuilder MIRBuilder(
I);
1792 Register OpReg =
I.getOperand(1).getReg();
1798 SPIRVTypeInst ExtType =
1805 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1809 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1812 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1815bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1816 SPIRVTypeInst ResType,
1819 unsigned Opcode)
const {
1820 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1823bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1824 SPIRVTypeInst ResType,
1827 unsigned Opcode)
const {
1828 MachineIRBuilder MIRBuilder(
I);
1837 SPIRVTypeInst WorkingType =
1844 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {SrcReg}, SPIRV::OpUConvert))
1848 if (!selectOpWithSrcs(LowCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1856 IsScalar ? SPIRV::OpShiftRightLogicalS : SPIRV::OpShiftRightLogicalV;
1858 if (!selectOpWithSrcs(Shift, SrcType,
I, {SrcReg, ShiftAmount}, ShiftOp))
1862 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {Shift}, SPIRV::OpUConvert))
1866 if (!selectOpWithSrcs(HighCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1871 if (!selectOpWithSrcs(Sum, WorkingType,
I, {HighCount, LowCount},
1872 IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV))
1876 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1877 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1880bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1881 SPIRVTypeInst ResType,
1883 unsigned Opcode)
const {
1888 if (!STI.getTargetTriple().isVulkanOS())
1889 return selectUnOp(ResVReg, ResType,
I, Opcode);
1891 Register OpReg =
I.getOperand(1).getReg();
1894 : SPIRV::OpUConvert;
1898 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1900 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1902 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1904 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1908bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1909 SPIRVTypeInst ResType,
1911 unsigned Opcode)
const {
1913 Register SrcReg =
I.getOperand(1).getReg();
1918 unsigned DefOpCode = DefIt->getOpcode();
1919 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1922 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1923 DefOpCode = VRD->getOpcode();
1925 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1926 DefOpCode == TargetOpcode::G_CONSTANT ||
1927 DefOpCode == SPIRV::OpVariable ||
1928 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1929 DefOpCode == SPIRV::OpConstantI) {
1935 uint32_t SpecOpcode = 0;
1937 case SPIRV::OpConvertPtrToU:
1938 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1940 case SPIRV::OpConvertUToPtr:
1941 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1946 TII.get(SPIRV::OpSpecConstantOp))
1956 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1960bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1961 SPIRVTypeInst ResType,
1962 MachineInstr &
I)
const {
1963 Register OpReg =
I.getOperand(1).getReg();
1964 SPIRVTypeInst OpType =
1967 return diagnoseUnsupported(
1968 I,
"incompatible result and operand types in a bitcast");
1969 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1980 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1981 if (
MemOp->isNonTemporal())
1982 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1984 if (!ST->isShader() &&
MemOp->getAlign().value())
1985 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1989 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1990 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1994 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1996 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
2000 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
2004 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
2006 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
2018 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2020 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2022 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2026bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2027 SPIRVTypeInst ResType,
2028 MachineInstr &
I)
const {
2030 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2035 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2036 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2038 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2040 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2044 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2048 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2049 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2050 I.getDebugLoc(),
I);
2054 MachineIRBuilder MIRBuilder(
I);
2056 if (
I.getNumMemOperands()) {
2057 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2058 if (MemOp->isAtomic())
2059 return selectAtomicLoad(ResVReg, ResType,
I);
2062 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2066 if (!
I.getNumMemOperands()) {
2067 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2069 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2078Register SPIRVInstructionSelector::createPtrSizedIntReg(
2079 MachineIRBuilder &MIRBuilder)
const {
2080 SPIRVTypeInst IntType =
2090SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2091 MachineIRBuilder &MIRBuilder)
const {
2092 SPIRVTypeInst IntType =
2094 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2095 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2103Register SPIRVInstructionSelector::castPtrToPtrToInt(
2104 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2105 MachineIRBuilder &MIRBuilder)
const {
2106 SPIRVTypeInst IntType =
2108 SPIRVTypeInst PtrType =
2122bool SPIRVInstructionSelector::selectAtomicPtrValue(
2123 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2124 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2132 Register IntResult = EmitAtomic(IntType);
2134 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2142bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2143 SPIRVTypeInst ResType,
2144 MachineInstr &
I)
const {
2145 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2148 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2151 return diagnoseUnsupported(
2152 I,
"Lowering to SPIR-V of atomic load is only "
2153 "allowed for integer, floating point or pointer types");
2155 assert(
I.getNumMemOperands());
2156 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2157 assert(MemOp.isAtomic());
2159 uint32_t
Scope =
static_cast<uint32_t
>(
2161 Register ScopeReg = buildI32Constant(Scope,
I);
2167 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2168 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2171 Register MemSemReg = buildI32Constant(Sem,
I);
2173 MachineIRBuilder MIRBuilder(
I);
2177 return diagnoseUnsupported(
2178 I,
"Lowering to SPIR-V of atomic load is only "
2179 "allowed for pointer types for physical addressing model");
2184 SPIRV::StorageClass::StorageClass SC =
2186 return selectAtomicPtrValue(
2187 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2188 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2189 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2200 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2211bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2213 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2214 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2219 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2220 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2222 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2227 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2231 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2232 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2233 SPIRVTypeInst SampledType =
2235 SPIRVTypeInst StoreValCompType =
2237 if (StoreValCompType && StoreValCompType != SampledType) {
2240 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2243 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2248 StoreVal = PackedReg;
2251 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2252 TII.get(SPIRV::OpImageWrite))
2258 if (sampledTypeIsSignedInteger(LLVMHandleType))
2261 BMI.constrainAllUses(
TII,
TRI, RBI);
2268 if (PointeeTy && PointeeTy->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
2269 StoreTy->
getOpcode() != SPIRV::OpTypeVectorIdEXT &&
2271 MachineInstr *StoreValDef =
getVRegDef(*MRI, StoreVal);
2283 if (
I.getNumMemOperands()) {
2284 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2285 if (MemOp->isAtomic())
2286 return selectAtomicStore(
I);
2293 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2294 PtrSC == SPIRV::StorageClass::Input ||
2295 PtrSC == SPIRV::StorageClass::PushConstant)
2296 return diagnoseUnsupported(
2297 I,
"store into a read-only SPIR-V storage class is not allowed");
2299 MachineIRBuilder MIRBuilder(
I);
2301 if (!
I.getNumMemOperands()) {
2302 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2304 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2313bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2314 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2317 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2318 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2323 if (!PointeeType && PtrType &&
2324 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2327 return diagnoseUnsupported(
I,
2328 "Lowering to SPIR-V of atomic store is only "
2329 "allowed for integer or floating point types");
2331 assert(
I.getNumMemOperands());
2332 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2333 assert(MemOp.isAtomic());
2335 uint32_t
Scope =
static_cast<uint32_t
>(
2337 Register ScopeReg = buildI32Constant(Scope,
I);
2343 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2344 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2347 Register MemSemReg = buildI32Constant(Sem,
I);
2348 MachineIRBuilder MIRBuilder(
I);
2352 return diagnoseUnsupported(
2353 I,
"Lowering to SPIR-V of atomic store is only "
2354 "allowed for pointer types for physical addressing model");
2359 SPIRV::StorageClass::StorageClass SC =
2361 return selectAtomicPtrValue(
2362 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2364 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2377 return diagnoseUnsupported(
I,
2378 "Lowering to SPIR-V of atomic store is only "
2379 "allowed for integer or floating point types");
2381 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2391bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2392 SPIRVTypeInst ResType,
2393 MachineInstr &
I)
const {
2394 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2402 const Register PtrsReg =
I.getOperand(2).getReg();
2403 const uint32_t
Alignment =
I.getOperand(3).getImm();
2404 const Register MaskReg =
I.getOperand(4).getReg();
2405 const Register PassthruReg =
I.getOperand(5).getReg();
2406 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2410 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2421bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2422 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2429 const Register ValuesReg =
I.getOperand(1).getReg();
2430 const Register PtrsReg =
I.getOperand(2).getReg();
2431 const uint32_t
Alignment =
I.getOperand(3).getImm();
2432 const Register MaskReg =
I.getOperand(4).getReg();
2433 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2437 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2446bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2447 const Twine &
Msg)
const {
2448 const Function &
F =
I.getMF()->getFunction();
2449 F.getContext().diagnose(
2450 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2454bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2455 SPIRVTypeInst ResType,
2456 MachineInstr &
I)
const {
2457 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2458 return diagnoseUnsupported(
2459 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2460 "SPIR-V extension: SPV_INTEL_variable_length_array");
2462 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2469bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2470 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2471 return diagnoseUnsupported(
2473 "llvm.stackrestore intrinsic: this instruction requires the following "
2474 "SPIR-V extension: SPV_INTEL_variable_length_array");
2475 if (!
I.getOperand(0).isReg())
2478 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2479 .
addUse(
I.getOperand(0).getReg())
2485SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2486 MachineIRBuilder MIRBuilder(
I);
2487 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2494 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2498 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2499 Type *ArrTy = ArrayType::get(ValTy, Num);
2501 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2504 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2515 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2516 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2517 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2518 : SPIRV::OpVariable))
2521 .
addImm(SPIRV::StorageClass::UniformConstant);
2534bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2537 Register DstReg =
I.getOperand(0).getReg();
2541 return diagnoseUnsupported(
2542 I,
"OpCopyMemory requires operands to have the same type");
2547 return diagnoseUnsupported(
2548 I,
"Unable to determine pointee type size for OpCopyMemory");
2549 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2550 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2551 return diagnoseUnsupported(
2552 I,
"OpCopyMemory requires the size to match the pointee type size");
2553 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2556 if (
I.getNumMemOperands()) {
2557 MachineIRBuilder MIRBuilder(
I);
2564bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2567 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2568 .
addUse(
I.getOperand(0).getReg())
2570 .
addUse(
I.getOperand(2).getReg());
2571 if (
I.getNumMemOperands()) {
2572 MachineIRBuilder MIRBuilder(
I);
2579bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2580 MachineInstr &
I)
const {
2582 Register SizeReg =
I.getOperand(2).getReg();
2584 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2588 Register SrcReg =
I.getOperand(1).getReg();
2589 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2590 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2591 Register VarReg = getOrCreateMemSetGlobal(
I);
2594 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2596 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2598 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2602 if (!selectCopyMemory(
I, SrcReg))
2605 if (!selectCopyMemorySized(
I, SrcReg))
2608 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2609 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2614bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2615 SPIRVTypeInst ResType,
2618 unsigned NegateOpcode)
const {
2620 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2621 uint32_t
Scope =
static_cast<uint32_t
>(
2623 MemOp->getSyncScopeID()));
2624 Register ScopeReg = buildI32Constant(Scope,
I);
2626 Register Ptr =
I.getOperand(1).getReg();
2627 uint32_t ScSem =
static_cast<uint32_t
>(
2631 Register MemSemReg = buildI32Constant(
2635 Register ValueReg =
I.getOperand(2).getReg();
2636 if (NegateOpcode != 0) {
2639 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2645 if (NewOpcode != SPIRV::OpAtomicExchange)
2646 return diagnoseUnsupported(
2647 I,
"Lowering to SPIR-V of this atomic operation is not "
2648 "allowed for pointer types");
2650 return diagnoseUnsupported(
2651 I,
"Lowering to SPIR-V of atomic exchange is only "
2652 "allowed for pointer types for physical addressing model");
2659 MachineIRBuilder MIRBuilder(
I);
2661 return selectAtomicPtrValue(
2662 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2664 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2665 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2666 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2674 return ExchangeResReg;
2678 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2689bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2690 unsigned ArgI =
I.getNumOperands() - 1;
2692 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2693 SPIRVTypeInst SrcType =
2697 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2701 unsigned CurrentIndex = 0;
2702 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2703 Register ResVReg =
I.getOperand(i).getReg();
2706 LLT ResLLT = MRI->
getType(ResVReg);
2712 ResType = ScalarType;
2721 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2727 for (
unsigned j = 0;
j < NumElements; ++
j) {
2728 MIB.
addImm(CurrentIndex + j);
2730 CurrentIndex += NumElements;
2734 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2746bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2749 ? SPIRV::MemorySemantics::UniformMemory |
2750 SPIRV::MemorySemantics::WorkgroupMemory |
2751 SPIRV::MemorySemantics::ImageMemory
2752 : SPIRV::MemorySemantics::WorkgroupMemory |
2753 SPIRV::MemorySemantics::CrossWorkgroupMemory |
2754 SPIRV::MemorySemantics::ImageMemory;
2756 STI.getTargetTriple(),
static_cast<uint32_t
>(
getMemSemantics(AO)), ScSem);
2757 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
2761 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2763 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2770bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2771 SPIRVTypeInst ResType,
2773 unsigned Opcode)
const {
2774 Type *ResTy =
nullptr;
2777 return diagnoseUnsupported(
2779 "Not enough info to select the arithmetic with overflow instruction");
2781 return diagnoseUnsupported(
I,
2782 "Expect struct type result for the arithmetic "
2783 "with overflow instruction");
2789 MachineIRBuilder MIRBuilder(
I);
2791 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2792 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2799 Register ZeroReg = buildZerosVal(ResType,
I);
2804 if (ResName.
size() > 0)
2812 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2813 MIB.
addUse(
I.getOperand(i).getReg());
2818 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2819 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2821 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2822 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2829 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2830 .
addDef(
I.getOperand(1).getReg())
2838bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2839 SPIRVTypeInst ResType,
2840 MachineInstr &
I)
const {
2842 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2843 Register Ptr =
I.getOperand(2).getReg();
2844 Register ScopeReg =
I.getOperand(5).getReg();
2845 Register MemSemEqReg =
I.getOperand(6).getReg();
2846 Register MemSemNeqReg =
I.getOperand(7).getReg();
2848 Register Val =
I.getOperand(4).getReg();
2852 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2871 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2878 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2890 case SPIRV::StorageClass::DeviceOnlyINTEL:
2891 case SPIRV::StorageClass::HostOnlyINTEL:
2900 bool IsGRef =
false;
2901 bool IsAllowedRefs =
2903 unsigned Opcode = It.getOpcode();
2904 if (Opcode == SPIRV::OpConstantComposite ||
2905 Opcode == SPIRV::OpSpecConstantComposite ||
2906 Opcode == SPIRV::OpVariable ||
2907 Opcode == SPIRV::OpUntypedVariableKHR ||
2908 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2909 return IsGRef = true;
2910 return Opcode == SPIRV::OpName;
2912 return IsAllowedRefs && IsGRef;
2915Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2916 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2918 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2922SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2924 uint32_t Opcode)
const {
2925 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2926 TII.get(SPIRV::OpSpecConstantOp))
2934SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2935 SPIRVTypeInst SrcPtrTy)
const {
2936 SPIRVTypeInst GenericPtrTy =
2940 SPIRV::StorageClass::Generic),
2944 MachineInstrBuilder MIB = buildSpecConstantOp(
2946 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2956bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2957 SPIRVTypeInst ResType,
2958 MachineInstr &
I)
const {
2962 Register SrcPtr =
I.getOperand(1).getReg();
2967 return BuildCOPY(ResVReg, SrcPtr,
I);
2977 unsigned SpecOpcode = [&]() ->
unsigned {
2978 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2979 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2981 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2983 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
2991 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2993 .constrainAllUses(
TII,
TRI, RBI);
2995 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2997 buildSpecConstantOp(
2999 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
3000 .constrainAllUses(
TII,
TRI, RBI);
3007 return BuildCOPY(ResVReg, SrcPtr,
I);
3009 if ((SrcSC == SPIRV::StorageClass::Function &&
3010 DstSC == SPIRV::StorageClass::Private) ||
3011 (DstSC == SPIRV::StorageClass::Function &&
3012 SrcSC == SPIRV::StorageClass::Private))
3013 return BuildCOPY(ResVReg, SrcPtr,
I);
3017 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3020 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3023 SPIRVTypeInst GenericPtrTy =
3042 return selectUnOp(ResVReg, ResType,
I,
3043 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
3045 return selectUnOp(ResVReg, ResType,
I,
3046 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
3048 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3050 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3060bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3061 SPIRVTypeInst ResType,
3062 MachineInstr &
I)
const {
3064 return diagnoseUnsupported(
3065 I,
"G_PTRMASK is not supported with logical SPIR-V");
3070 Register PtrReg =
I.getOperand(1).getReg();
3071 Register MaskReg =
I.getOperand(2).getReg();
3090 ? SPIRV::OpBitwiseAndV
3091 : SPIRV::OpBitwiseAndS;
3114 return SPIRV::OpFOrdEqual;
3116 return SPIRV::OpFOrdGreaterThanEqual;
3118 return SPIRV::OpFOrdGreaterThan;
3120 return SPIRV::OpFOrdLessThanEqual;
3122 return SPIRV::OpFOrdLessThan;
3124 return SPIRV::OpFOrdNotEqual;
3126 return SPIRV::OpOrdered;
3128 return SPIRV::OpFUnordEqual;
3130 return SPIRV::OpFUnordGreaterThanEqual;
3132 return SPIRV::OpFUnordGreaterThan;
3134 return SPIRV::OpFUnordLessThanEqual;
3136 return SPIRV::OpFUnordLessThan;
3138 return SPIRV::OpFUnordNotEqual;
3140 return SPIRV::OpUnordered;
3150 return SPIRV::OpIEqual;
3152 return SPIRV::OpINotEqual;
3154 return SPIRV::OpSGreaterThanEqual;
3156 return SPIRV::OpSGreaterThan;
3158 return SPIRV::OpSLessThanEqual;
3160 return SPIRV::OpSLessThan;
3162 return SPIRV::OpUGreaterThanEqual;
3164 return SPIRV::OpUGreaterThan;
3166 return SPIRV::OpULessThanEqual;
3168 return SPIRV::OpULessThan;
3177 return SPIRV::OpPtrEqual;
3179 return SPIRV::OpPtrNotEqual;
3190 return SPIRV::OpLogicalEqual;
3192 return SPIRV::OpLogicalNotEqual;
3230bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3231 SPIRVTypeInst ResType,
3233 unsigned OpAnyOrAll)
const {
3234 assert(
I.getNumOperands() == 3);
3235 assert(
I.getOperand(2).isReg());
3237 Register InputRegister =
I.getOperand(2).getReg();
3240 assert(InputType &&
"VReg has no type assigned");
3244 assert(ResVReg ==
I.getOperand(0).getReg());
3245 return BuildCOPY(ResVReg, InputRegister,
I);
3249 unsigned SpirvNotEqualId =
3250 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3252 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3257 IsBoolTy ? InputRegister
3265 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3267 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3284bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3285 SPIRVTypeInst ResType,
3286 MachineInstr &
I)
const {
3287 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3290bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3291 SPIRVTypeInst ResType,
3292 MachineInstr &
I)
const {
3293 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3297bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3298 SPIRVTypeInst ResType,
3299 MachineInstr &
I)
const {
3300 assert(
I.getNumOperands() == 4);
3301 assert(
I.getOperand(2).isReg());
3302 assert(
I.getOperand(3).isReg());
3304 [[maybe_unused]] SPIRVTypeInst VecType =
3309 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3310 "dot product requires either a vector of at least 2 components or"
3311 " the SPV_EXT_long vector extension.");
3313 [[maybe_unused]] SPIRVTypeInst EltType =
3322 .
addUse(
I.getOperand(2).getReg())
3323 .
addUse(
I.getOperand(3).getReg())
3328bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3329 SPIRVTypeInst ResType,
3332 assert(
I.getNumOperands() == 4);
3333 assert(
I.getOperand(2).isReg());
3334 assert(
I.getOperand(3).isReg());
3337 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3341 .
addUse(
I.getOperand(2).getReg())
3342 .
addUse(
I.getOperand(3).getReg())
3349bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3350 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3351 assert(
I.getNumOperands() == 4);
3352 assert(
I.getOperand(2).isReg());
3353 assert(
I.getOperand(3).isReg());
3357 Register Vec0 =
I.getOperand(2).getReg();
3358 Register Vec1 =
I.getOperand(3).getReg();
3362 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3371 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3372 "dot product requires either a vector of at least 2 components "
3373 "or the SPV_EXT_long_vector extension.");
3376 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3386 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3397 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3409bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3410 SPIRVTypeInst ResType,
3411 MachineInstr &
I)
const {
3413 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3416 .
addUse(
I.getOperand(2).getReg())
3421bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3422 SPIRVTypeInst ResType,
3423 MachineInstr &
I)
const {
3425 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3428 .
addUse(
I.getOperand(2).getReg())
3433bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3434 SPIRVTypeInst ResType,
3435 MachineInstr &
I)
const {
3437 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3440 .
addUse(
I.getOperand(2).getReg())
3445bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3446 SPIRVTypeInst ResType,
3447 MachineInstr &
I)
const {
3449 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3452 .
addUse(
I.getOperand(2).getReg())
3457template <
bool Signed>
3458bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3459 SPIRVTypeInst ResType,
3460 MachineInstr &
I)
const {
3461 assert(
I.getNumOperands() == 5);
3462 assert(
I.getOperand(2).isReg());
3463 assert(
I.getOperand(3).isReg());
3464 assert(
I.getOperand(4).isReg());
3467 Register Acc =
I.getOperand(2).getReg();
3471 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3473 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3478 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3481 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3493template <
bool Signed>
3494bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3495 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3496 assert(
I.getNumOperands() == 5);
3497 assert(
I.getOperand(2).isReg());
3498 assert(
I.getOperand(3).isReg());
3499 assert(
I.getOperand(4).isReg());
3502 Register Acc =
I.getOperand(2).getReg();
3508 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3512 for (
unsigned i = 0; i < 4; i++) {
3535 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3555 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3570bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3571 SPIRVTypeInst ResType,
3572 MachineInstr &
I)
const {
3573 assert(
I.getNumOperands() == 3);
3574 assert(
I.getOperand(2).isReg());
3576 Register VZero = buildZerosValF(ResType,
I);
3577 Register VOne = buildOnesValF(ResType,
I);
3579 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3582 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3584 .
addUse(
I.getOperand(2).getReg())
3591bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3592 SPIRVTypeInst ResType,
3593 MachineInstr &
I)
const {
3594 assert(
I.getNumOperands() == 3);
3595 assert(
I.getOperand(2).isReg());
3597 Register InputRegister =
I.getOperand(2).getReg();
3599 auto &
DL =
I.getDebugLoc();
3602 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3609 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3611 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3619 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3624 if (NeedsConversion) {
3625 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3636bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3637 SPIRVTypeInst ResType,
3639 unsigned Opcode)
const {
3643 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3649 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3650 BMI.addUse(
I.getOperand(J).getReg());
3657bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3660 bool WithGroupSync)
const {
3662 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3664 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3666 assert(((Scope != SPIRV::Scope::Workgroup) ||
3667 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3668 "Workgroup Scope must set WorkGroupMemory semantic "
3669 "in Barrier instruction");
3671 assert(((Scope != SPIRV::Scope::Device) ||
3672 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3673 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3674 "Device Scope must set UniformMemory and ImageMemory semantic "
3675 "in Barrier instruction");
3681 if (WithGroupSync) {
3682 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3686 Register ScopeReg = buildI32Constant(Scope,
I);
3687 Register MemSemReg = buildI32Constant(MemSem,
I);
3689 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3693bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3694 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3699 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3700 SPIRV::OpGroupNonUniformBallot))
3705 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3710 .
addImm(SPIRV::GroupOperation::Reduce)
3717bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3718 SPIRVTypeInst ResType,
3719 MachineInstr &
I)
const {
3724 Register InputReg =
I.getOperand(2).getReg();
3729 bool IsVector = NumElems > 1 ||
3730 (InputType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
3744 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3745 SPIRV::OpGroupNonUniformAllEqual);
3750 ElementResults.
reserve(NumElems);
3752 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3765 ElemInput = Extracted;
3771 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3782 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3793bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3794 SPIRVTypeInst ResType,
3795 MachineInstr &
I)
const {
3797 assert(
I.getNumOperands() == 3);
3799 auto Op =
I.getOperand(2);
3809 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3811 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3812 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3833 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3837 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3844bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3845 SPIRVTypeInst ResType,
3847 bool IsUnsigned)
const {
3848 return selectWaveReduce(
3849 ResVReg, ResType,
I, IsUnsigned,
3850 [&](
Register InputRegister,
bool IsUnsigned) {
3851 const bool IsFloatTy =
3853 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3854 : SPIRV::OpGroupNonUniformSMax;
3855 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3859bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3860 SPIRVTypeInst ResType,
3862 bool IsUnsigned)
const {
3863 return selectWaveReduce(
3864 ResVReg, ResType,
I, IsUnsigned,
3865 [&](
Register InputRegister,
bool IsUnsigned) {
3866 const bool IsFloatTy =
3868 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3869 : SPIRV::OpGroupNonUniformSMin;
3870 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3874bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3875 SPIRVTypeInst ResType,
3876 MachineInstr &
I)
const {
3877 return selectWaveReduce(ResVReg, ResType,
I,
false,
3878 [&](
Register InputRegister,
bool IsUnsigned) {
3880 InputRegister, SPIRV::OpTypeFloat);
3881 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3882 : SPIRV::OpGroupNonUniformIAdd;
3886bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3887 SPIRVTypeInst ResType,
3888 MachineInstr &
I)
const {
3889 return selectWaveReduce(ResVReg, ResType,
I,
false,
3890 [&](
Register InputRegister,
bool IsUnsigned) {
3892 InputRegister, SPIRV::OpTypeFloat);
3893 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3894 : SPIRV::OpGroupNonUniformIMul;
3898template <
typename PickOpcodeFn>
3899bool SPIRVInstructionSelector::selectWaveReduce(
3900 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3901 PickOpcodeFn &&PickOpcode)
const {
3902 assert(
I.getNumOperands() == 3);
3903 assert(
I.getOperand(2).isReg());
3905 Register InputRegister =
I.getOperand(2).getReg();
3909 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3912 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3918 .
addImm(SPIRV::GroupOperation::Reduce)
3919 .
addUse(
I.getOperand(2).getReg())
3924bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3925 SPIRVTypeInst ResType,
3927 unsigned Opcode)
const {
3928 return selectWaveReduce(
3929 ResVReg, ResType,
I,
false,
3930 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3933bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3934 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3935 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3936 [&](
Register InputRegister,
bool IsUnsigned) {
3938 InputRegister, SPIRV::OpTypeFloat);
3940 ? SPIRV::OpGroupNonUniformFAdd
3941 : SPIRV::OpGroupNonUniformIAdd;
3945bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3946 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3947 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3948 [&](
Register InputRegister,
bool IsUnsigned) {
3950 InputRegister, SPIRV::OpTypeFloat);
3952 ? SPIRV::OpGroupNonUniformFMul
3953 : SPIRV::OpGroupNonUniformIMul;
3957template <
typename PickOpcodeFn>
3958bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3959 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3960 PickOpcodeFn &&PickOpcode)
const {
3961 assert(
I.getNumOperands() == 3);
3962 assert(
I.getOperand(2).isReg());
3964 Register InputRegister =
I.getOperand(2).getReg();
3968 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3971 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3977 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3978 .
addUse(
I.getOperand(2).getReg())
3983bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3984 SPIRVTypeInst ResType,
3987 assert(
I.getNumOperands() == 3);
3988 assert(
I.getOperand(2).isReg());
3990 Register InputRegister =
I.getOperand(2).getReg();
3996 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
4007bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
4008 SPIRVTypeInst ResType,
4015 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
4020 : SPIRV::OpUConvert;
4022 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4025 ShiftOp = SPIRV::OpShiftRightLogicalV;
4030 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4031 TII.get(SPIRV::OpConstantComposite))
4034 for (
unsigned It = 0; It <
N; ++It)
4038 ShiftConst = CompositeReg;
4043 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
4048 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
4053 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
4058 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
4061bool SPIRVInstructionSelector::handle64BitOverflow(
4063 unsigned int Opcode,
4070 "handle64BitOverflow should only be used for integer types");
4072 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4074 MachineIRBuilder MIRBuilder(
I);
4076 SPIRVTypeInst I64x2Type =
4078 SPIRVTypeInst Vec2ResType =
4081 std::vector<Register> PartialRegs;
4083 unsigned CurrentComponent = 0;
4084 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4088 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4089 TII.get(SPIRV::OpVectorShuffle))
4094 .
addImm(CurrentComponent)
4095 .
addImm(CurrentComponent + 1);
4105 PartialRegs.push_back(SubVecReg);
4108 if (CurrentComponent != ComponentCount) {
4114 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4115 SPIRV::OpVectorExtractDynamic))
4124 PartialRegs.push_back(FinalElemResReg);
4128 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4129 SPIRV::OpCompositeConstruct);
4132bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4133 SPIRVTypeInst ResType,
4137 if (ComponentCount > 2)
4138 return handle64BitOverflow(
4139 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4141 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4143 MachineIRBuilder MIRBuilder(
I);
4147 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4151 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4156 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4163 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4164 TII.get(SPIRV::OpVectorShuffle))
4169 for (
unsigned J = 0; J < ComponentCount; ++J) {
4176 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4179bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4180 SPIRVTypeInst ResType,
4184 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4192bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4193 SPIRVTypeInst ResType,
4194 MachineInstr &
I)
const {
4195 Register OpReg =
I.getOperand(1).getReg();
4204 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4206 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4208 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4210 return SPIRVInstructionSelector::diagnoseUnsupported(
4211 I,
"G_BITREVERSE only support 16,32,64 bits.");
4215 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4226 unsigned AndOp = SPIRV::OpBitwiseAndS;
4227 unsigned OrOp = SPIRV::OpBitwiseOrS;
4228 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4229 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4230 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4232 AndOp = SPIRV::OpBitwiseAndV;
4233 OrOp = SPIRV::OpBitwiseOrV;
4234 ShlOp = SPIRV::OpShiftLeftLogicalV;
4235 ShrOp = SPIRV::OpShiftRightLogicalV;
4241 const unsigned Shift) ->
Register {
4244 (ResType->
getOpcode() != SPIRV::OpTypeVectorIdEXT ||
4251 Register MaskReg = CreateConst(Mask);
4252 Register ShiftReg = CreateConst(Shift);
4259 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4260 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4261 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4262 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4263 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4272 while ((Shift >>= 1) > 0) {
4279 return BuildCOPY(ResVReg, Result,
I);
4282bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4283 SPIRVTypeInst ResType,
4284 MachineInstr &
I)
const {
4285 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4286 "G_FREEZE must define and use a register");
4287 Register OpReg =
I.getOperand(1).getReg();
4291 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4304 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4305 if (
Def->getOpcode() == TargetOpcode::COPY)
4308 switch (
Def->getOpcode()) {
4309 case SPIRV::ASSIGN_TYPE:
4310 if (MachineInstr *AssignToDef =
4312 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4313 Reg =
Def->getOperand(2).getReg();
4316 case SPIRV::OpUndef:
4317 Reg =
Def->getOperand(1).getReg();
4320 unsigned DestOpCode;
4322 DestOpCode = SPIRV::OpConstantNull;
4323 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4324 "static undef/poison lowered to OpConstantNull\n");
4326 DestOpCode = TargetOpcode::COPY;
4328 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4329 "skipped, lowered as a copy of the operand\n");
4331 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4332 .
addDef(
I.getOperand(0).getReg())
4340bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4341 SPIRVTypeInst ResType,
4342 MachineInstr &
I)
const {
4346 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4350 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4355 for (
unsigned i =
I.getNumExplicitDefs();
4356 i <
I.getNumExplicitOperands() && IsConst; ++i)
4361 return diagnoseUnsupported(
4362 I,
"There must be at least two constituent operands in a vector");
4367 for (
unsigned i =
I.getNumExplicitDefs();
4368 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4369 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4374 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4381 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4382 TII.get(IsConst ? SPIRV::OpConstantComposite
4383 : SPIRV::OpCompositeConstruct))
4386 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4387 MIB.
addUse(
I.getOperand(i).getReg());
4392bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4393 SPIRVTypeInst ResType,
4394 MachineInstr &
I)
const {
4398 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4403 unsigned OpIdx =
I.getNumExplicitDefs();
4404 if (!
I.getOperand(OpIdx).isReg())
4408 Register OpReg =
I.getOperand(OpIdx).getReg();
4412 return diagnoseUnsupported(
4413 I,
"There must be at least two constituent operands in a vector");
4416 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4417 TII.get(IsConst ? SPIRV::OpConstantComposite
4418 : SPIRV::OpCompositeConstruct))
4421 for (
unsigned i = 0; i <
N; ++i)
4427bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4428 SPIRVTypeInst ResType,
4429 MachineInstr &
I)
const {
4435 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4437 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4438 TII.get(SPIRV::OpCompositeConstruct))
4441 for (
unsigned OpIdx =
I.getNumExplicitDefs();
4442 OpIdx <
I.getNumExplicitOperands(); ++OpIdx)
4443 MIB.
addUse(
I.getOperand(OpIdx).getReg());
4448bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4449 SPIRVTypeInst ResType,
4450 MachineInstr &
I)
const {
4456 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4458 Opcode = SPIRV::OpDemoteToHelperInvocation;
4460 Opcode = SPIRV::OpKill;
4465 ToErase.eraseFromParent();
4474bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4475 SPIRVTypeInst ResType,
unsigned CmpOpc,
4476 MachineInstr &
I)
const {
4477 Register Cmp0 =
I.getOperand(2).getReg();
4478 Register Cmp1 =
I.getOperand(3).getReg();
4481 "CMP operands should have the same type");
4482 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4492bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4493 SPIRVTypeInst ResType,
4494 MachineInstr &
I)
const {
4495 auto Pred =
I.getOperand(1).getPredicate();
4498 Register CmpOperand =
I.getOperand(2).getReg();
4500 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4505 Register Op1 =
I.getOperand(3).getReg();
4509 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4514 I.getOperand(3).setReg(NewOp1);
4520 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4524SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4525 SPIRVTypeInst ResType)
const {
4527 SPIRVTypeInst SpvI32Ty =
4530 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4537 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4540 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4543 .
addImm(APInt(32, Val).getZExtValue());
4545 GR.
add(ConstInt,
MI);
4552Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4553 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4555 SPIRVTypeInst SpvI32Ty =
4557 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4562 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4563 MachineInstr *
MI =
nullptr;
4567 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4571 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4572 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4578 GR.
add(ConstInt,
MI);
4583bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4584 SPIRVTypeInst ResType,
4585 MachineInstr &
I)
const {
4587 return selectCmp(ResVReg, ResType, CmpOp,
I);
4590bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4591 SPIRVTypeInst ResType,
4592 MachineInstr &
I)
const {
4594 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4604 MachineIRBuilder MIRBuilder(
I);
4611 APFloat ConstVal(3.3219280948873623);
4615 APFloat::rmNearestTiesToEven, &LosesInfo);
4620 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
4622 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4623 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4625 if (!selectExtInst(ResVReg, ResType,
I,
4626 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4636Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4637 MachineInstr &
I)
const {
4645bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4651 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4659 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4662 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4663 Def->getOpcode() == SPIRV::OpConstantI)
4676 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4677 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4679 Intrinsic::spv_const_composite)) {
4680 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4681 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4682 if (!IsZero(
Def->getOperand(i).getReg()))
4691Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4692 MachineInstr &
I)
const {
4701Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4702 MachineInstr &
I)
const {
4712 SPIRVTypeInst ResType,
4713 MachineInstr &
I)
const {
4722bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4723 SPIRVTypeInst ResType,
4724 MachineInstr &
I)
const {
4725 Register SelectFirstArg =
I.getOperand(2).getReg();
4726 Register SelectSecondArg =
I.getOperand(3).getReg();
4740 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4741 }
else if (IsPtrTy) {
4742 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4744 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4747 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4748 "boolean condition");
4750 Opcode = SPIRV::OpSelectSFSCond;
4751 }
else if (IsPtrTy) {
4752 Opcode = SPIRV::OpSelectSPSCond;
4754 Opcode = SPIRV::OpSelectSISCond;
4757 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4760 .
addUse(
I.getOperand(1).getReg())
4769bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4770 SPIRVTypeInst ResType,
4772 MachineInstr &InsertAt,
4773 bool IsSigned)
const {
4775 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4776 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4777 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4779 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4791bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4792 SPIRVTypeInst ResType,
4793 MachineInstr &
I,
bool IsSigned,
4794 unsigned Opcode)
const {
4795 Register SrcReg =
I.getOperand(1).getReg();
4806 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4808 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4811bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4812 SPIRVTypeInst ResType, MachineInstr &
I,
4813 bool IsSigned)
const {
4814 Register SrcReg =
I.getOperand(1).getReg();
4816 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4820 if (ResType == SrcType)
4821 return BuildCOPY(ResVReg, SrcReg,
I);
4823 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4824 return selectUnOp(ResVReg, ResType,
I, Opcode);
4827bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4828 SPIRVTypeInst ResType,
4830 bool IsSigned)
const {
4831 MachineIRBuilder MIRBuilder(
I);
4832 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
4837 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4845 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4848 .
addUse(
I.getOperand(1).getReg())
4849 .
addUse(
I.getOperand(2).getReg())
4854 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4857 .
addUse(
I.getOperand(1).getReg())
4858 .
addUse(
I.getOperand(2).getReg())
4866 unsigned SelectOpcode =
4867 (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4869 ? SPIRV::OpSelectVIVCond
4870 : SPIRV::OpSelectSISCond;
4875 .
addUse(buildOnesVal(
true, ResType,
I))
4876 .
addUse(buildZerosVal(ResType,
I))
4883 .
addUse(buildOnesVal(
false, ResType,
I))
4888bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4891 SPIRVTypeInst IntTy,
4892 SPIRVTypeInst BoolTy)
const {
4896 isVectorType(IntTy) ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4898 Register One = buildOnesVal(
false, IntTy,
I);
4906 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4915bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4916 SPIRVTypeInst ResType,
4917 MachineInstr &
I)
const {
4918 Register IntReg =
I.getOperand(1).getReg();
4921 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4922 if (ArgType == ResType)
4923 return BuildCOPY(ResVReg, IntReg,
I);
4925 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4926 return selectUnOp(ResVReg, ResType,
I, Opcode);
4929bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4930 SPIRVTypeInst ResType,
4931 MachineInstr &
I)
const {
4932 unsigned Opcode =
I.getOpcode();
4933 unsigned TpOpcode = ResType->
getOpcode();
4935 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
4936 assert(Opcode == TargetOpcode::G_CONSTANT &&
4937 I.getOperand(1).getCImm()->isZero());
4938 MachineBasicBlock &DepMBB =
I.getMF()->front();
4941 }
else if (TpOpcode == SPIRV::OpTypeVectorIdEXT) {
4946 "Expected <1 x T> Vector!");
4947 if (Opcode == TargetOpcode::G_FCONSTANT)
4953 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4961 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4964bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4965 SPIRVTypeInst ResType,
4966 MachineInstr &
I)
const {
4967 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4974bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4975 SPIRVTypeInst ResType,
4976 MachineInstr &
I)
const {
4978 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4982 .
addUse(
I.getOperand(3).getReg())
4984 .
addUse(
I.getOperand(2).getReg());
4985 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4991bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4992 SPIRVTypeInst ResType,
4993 MachineInstr &
I)
const {
4994 Type *MaybeResTy =
nullptr;
4999 "Expected aggregate type for extractv instruction");
5001 SPIRV::AccessQualifier::ReadWrite,
false);
5005 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5008 .
addUse(
I.getOperand(2).getReg());
5009 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
5015bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
5016 SPIRVTypeInst ResType,
5017 MachineInstr &
I)
const {
5018 if (
getImm(
I.getOperand(4), MRI))
5019 return selectInsertVal(ResVReg, ResType,
I);
5021 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
5024 .
addUse(
I.getOperand(2).getReg())
5025 .
addUse(
I.getOperand(3).getReg())
5026 .
addUse(
I.getOperand(4).getReg())
5031bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
5032 SPIRVTypeInst ResType,
5033 MachineInstr &
I)
const {
5034 if (
getImm(
I.getOperand(3), MRI))
5035 return selectExtractVal(ResVReg, ResType,
I);
5037 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
5040 .
addUse(
I.getOperand(2).getReg())
5041 .
addUse(
I.getOperand(3).getReg())
5046bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
5047 SPIRVTypeInst ResType,
5048 MachineInstr &
I)
const {
5049 const bool IsGEPInBounds =
I.getOperand(2).getImm();
5052 const bool UseUntypedPointers =
5053 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
5058 if (UseUntypedPointers) {
5060 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
5061 : SPIRV::OpUntypedAccessChainKHR;
5063 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
5064 : SPIRV::OpUntypedPtrAccessChainKHR;
5073 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
5075 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
5076 : SPIRV::OpPtrAccessChain;
5081 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5086 if (UseUntypedPointers) {
5101 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5102 Def->getOperand(1).isReg())
5104 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5105 if (
const auto *GVar =
5108 SPIRV::AccessQualifier::ReadWrite,
5112 return diagnoseUnsupported(
5113 I,
"could not deduce the base type of an untyped access chain");
5118 Res.addUse(BaseReg);
5120 const bool IsAccessChainOpcode =
5121 (Opcode == SPIRV::OpAccessChain ||
5122 Opcode == SPIRV::OpInBoundsAccessChain ||
5123 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5124 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5126 assert((!IsAccessChainOpcode || (
getImm(
I.getOperand(4), MRI) &&
5127 foldImm(
I.getOperand(4), MRI) == 0)) &&
5128 "Cannot translate GEP to OpAccessChain.");
5131 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5132 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5133 Res.addUse(
I.getOperand(i).getReg());
5134 Res.constrainAllUses(
TII,
TRI, RBI);
5143 if (Extract.
getOpcode() == SPIRV::OpCompositeExtract) {
5149 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5155 TII.get(SPIRV::OpCompositeInsert))
5169bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5171 unsigned Lim =
I.getNumExplicitOperands();
5172 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5173 Register OpReg =
I.getOperand(i).getReg();
5174 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
5176 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
5177 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5178 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5191 SPIRVTypeInst WrapType = OpType;
5192 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5194 SPIRV::StorageClass::CodeSectionINTEL) {
5196 SPIRV::StorageClass::Function,
I);
5203 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5204 TII.get(SPIRV::OpSpecConstantOp))
5207 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5209 GR.
add(OpDefine, MIB);
5215bool SPIRVInstructionSelector::selectDerivativeInst(
5216 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5217 const unsigned DPdOpCode)
const {
5220 if (!errorIfInstrOutsideShader(
I))
5226 Register SrcReg =
I.getOperand(2).getReg();
5231 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5234 .
addUse(
I.getOperand(2).getReg());
5236 MachineIRBuilder MIRBuilder(
I);
5239 if (componentCount != 1)
5247 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5252 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5257 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5265bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5266 SPIRVTypeInst ResType,
5267 MachineInstr &
I)
const {
5271 case Intrinsic::spv_load:
5272 return selectLoad(ResVReg, ResType,
I);
5273 case Intrinsic::spv_atomic_load:
5274 return selectAtomicLoad(ResVReg, ResType,
I);
5275 case Intrinsic::spv_store:
5276 return selectStore(
I);
5277 case Intrinsic::spv_atomic_store:
5278 return selectAtomicStore(
I);
5279 case Intrinsic::spv_extractv:
5280 return selectExtractVal(ResVReg, ResType,
I);
5281 case Intrinsic::spv_insertv:
5282 return selectInsertVal(ResVReg, ResType,
I);
5283 case Intrinsic::spv_extractelt:
5284 return selectExtractElt(ResVReg, ResType,
I);
5285 case Intrinsic::spv_insertelt:
5286 return selectInsertElt(ResVReg, ResType,
I);
5287 case Intrinsic::spv_gep:
5288 return selectGEP(ResVReg, ResType,
I);
5289 case Intrinsic::spv_bitcast: {
5290 Register OpReg =
I.getOperand(2).getReg();
5291 SPIRVTypeInst OpType =
5295 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5297 case Intrinsic::spv_unref_global:
5298 case Intrinsic::spv_init_global: {
5299 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
5304 Register GVarVReg =
MI->getOperand(0).getReg();
5305 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5310 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5312 MI->eraseFromParent();
5316 case Intrinsic::spv_undef: {
5317 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5323 case Intrinsic::spv_poison:
5324 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5329 case Intrinsic::spv_freeze:
5330 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5333 .
addUse(
I.getOperand(2).getReg())
5336 case Intrinsic::spv_named_boolean_spec_constant: {
5337 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5338 : SPIRV::OpSpecConstantFalse;
5340 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5341 .
addDef(
I.getOperand(0).getReg())
5344 unsigned SpecId =
I.getOperand(2).getImm();
5346 SPIRV::Decoration::SpecId, {SpecId});
5350 case Intrinsic::spv_const_composite: {
5352 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5358 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5360 std::function<bool(
Register)> HasSpecConstOperand =
5370 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5371 J < Def->getNumExplicitOperands(); ++J) {
5372 if (
Def->getOperand(J).isReg() &&
5373 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5379 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5380 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5381 : SPIRV::OpConstantComposite;
5382 unsigned ContinuedOpc = HasSpecConst
5383 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5384 : SPIRV::OpConstantCompositeContinuedINTEL;
5385 MachineIRBuilder MIR(
I);
5387 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5389 for (
auto *Instr : Instructions) {
5390 Instr->setDebugLoc(
I.getDebugLoc());
5395 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5402 case Intrinsic::spv_assign_name: {
5403 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5404 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5405 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5406 i <
I.getNumExplicitOperands(); ++i) {
5407 MIB.
addImm(
I.getOperand(i).getImm());
5412 case Intrinsic::spv_switch: {
5413 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5414 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5415 if (
I.getOperand(i).isReg())
5416 MIB.
addReg(
I.getOperand(i).getReg());
5417 else if (
I.getOperand(i).isCImm())
5418 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5419 else if (
I.getOperand(i).isMBB())
5420 MIB.
addMBB(
I.getOperand(i).getMBB());
5427 case Intrinsic::spv_loop_merge: {
5428 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5429 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5430 if (
I.getOperand(i).isMBB())
5431 MIB.
addMBB(
I.getOperand(i).getMBB());
5438 case Intrinsic::spv_loop_control_intel: {
5440 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5441 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5446 case Intrinsic::spv_selection_merge: {
5448 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5449 assert(
I.getOperand(1).isMBB() &&
5450 "operand 1 to spv_selection_merge must be a basic block");
5451 MIB.
addMBB(
I.getOperand(1).getMBB());
5452 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5456 case Intrinsic::spv_cmpxchg:
5457 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5458 case Intrinsic::spv_unreachable:
5459 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5462 case Intrinsic::spv_abort:
5463 return selectAbort(
I);
5464 case Intrinsic::spv_alloca:
5465 return selectFrameIndex(ResVReg, ResType,
I);
5466 case Intrinsic::spv_alloca_array:
5467 return selectAllocaArray(ResVReg, ResType,
I);
5468 case Intrinsic::spv_assume:
5470 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5471 .
addUse(
I.getOperand(1).getReg())
5476 case Intrinsic::spv_expect:
5478 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5481 .
addUse(
I.getOperand(2).getReg())
5482 .
addUse(
I.getOperand(3).getReg())
5487 case Intrinsic::arithmetic_fence:
5488 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5489 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5492 .
addUse(
I.getOperand(2).getReg())
5496 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5498 case Intrinsic::spv_thread_id:
5504 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5506 case Intrinsic::spv_thread_id_in_group:
5512 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5514 case Intrinsic::spv_group_id:
5520 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5522 case Intrinsic::spv_flattened_thread_id_in_group:
5529 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5531 case Intrinsic::spv_workgroup_size:
5532 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5534 case Intrinsic::spv_global_size:
5535 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5537 case Intrinsic::spv_global_offset:
5538 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5540 case Intrinsic::spv_num_workgroups:
5541 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5543 case Intrinsic::spv_subgroup_size:
5544 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5546 case Intrinsic::spv_num_subgroups:
5547 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5549 case Intrinsic::spv_subgroup_id:
5550 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5551 case Intrinsic::spv_subgroup_local_invocation_id:
5552 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5553 ResVReg, ResType,
I);
5554 case Intrinsic::spv_subgroup_max_size:
5555 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5557 case Intrinsic::spv_fdot:
5558 return selectFloatDot(ResVReg, ResType,
I);
5559 case Intrinsic::spv_udot:
5560 case Intrinsic::spv_sdot:
5561 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5563 return selectIntegerDot(ResVReg, ResType,
I,
5564 IID == Intrinsic::spv_sdot);
5565 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5566 case Intrinsic::spv_dot4add_i8packed:
5567 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5569 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5570 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5571 case Intrinsic::spv_dot4add_u8packed:
5572 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5574 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5575 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5576 case Intrinsic::spv_all:
5577 return selectAll(ResVReg, ResType,
I);
5578 case Intrinsic::spv_any:
5579 return selectAny(ResVReg, ResType,
I);
5580 case Intrinsic::spv_distance:
5581 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5582 case Intrinsic::spv_lerp:
5583 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5584 case Intrinsic::spv_length:
5585 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5586 case Intrinsic::spv_degrees:
5587 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5588 case Intrinsic::spv_faceforward:
5589 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5590 case Intrinsic::spv_frac:
5591 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5592 case Intrinsic::spv_isinf:
5593 return selectOpIsInf(ResVReg, ResType,
I);
5594 case Intrinsic::spv_isnan:
5595 return selectOpIsNan(ResVReg, ResType,
I);
5596 case Intrinsic::spv_isfinite:
5597 return selectOpIsFinite(ResVReg, ResType,
I);
5598 case Intrinsic::spv_isnormal:
5599 return selectOpIsNormal(ResVReg, ResType,
I);
5600 case Intrinsic::spv_normalize:
5601 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5602 case Intrinsic::spv_refract:
5603 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5604 case Intrinsic::spv_reflect:
5605 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5606 case Intrinsic::spv_rsqrt:
5607 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5608 case Intrinsic::spv_sign:
5609 return selectSign(ResVReg, ResType,
I);
5610 case Intrinsic::spv_smoothstep:
5611 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5612 case Intrinsic::spv_firstbituhigh:
5613 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5614 case Intrinsic::spv_firstbitshigh:
5615 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5616 case Intrinsic::spv_firstbitlow:
5617 return selectFirstBitLow(ResVReg, ResType,
I);
5618 case Intrinsic::spv_all_memory_barrier:
5619 return selectBarrierInst(
I, SPIRV::Scope::Device,
5620 SPIRV::MemorySemantics::UniformMemory |
5621 SPIRV::MemorySemantics::ImageMemory |
5622 SPIRV::MemorySemantics::WorkgroupMemory,
5624 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5625 return selectBarrierInst(
I, SPIRV::Scope::Device,
5626 SPIRV::MemorySemantics::UniformMemory |
5627 SPIRV::MemorySemantics::ImageMemory |
5628 SPIRV::MemorySemantics::WorkgroupMemory,
5630 case Intrinsic::spv_device_memory_barrier:
5631 return selectBarrierInst(
I, SPIRV::Scope::Device,
5632 SPIRV::MemorySemantics::UniformMemory |
5633 SPIRV::MemorySemantics::ImageMemory,
5635 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5636 return selectBarrierInst(
I, SPIRV::Scope::Device,
5637 SPIRV::MemorySemantics::UniformMemory |
5638 SPIRV::MemorySemantics::ImageMemory,
5640 case Intrinsic::spv_group_memory_barrier:
5641 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5642 SPIRV::MemorySemantics::WorkgroupMemory,
5644 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5645 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5646 SPIRV::MemorySemantics::WorkgroupMemory,
5648 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5649 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5650 SPIRV::StorageClass::StorageClass ResSC =
5653 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5654 "from the Generic storage class");
5655 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5663 case Intrinsic::spv_lifetime_start:
5664 case Intrinsic::spv_lifetime_end: {
5665 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5666 : SPIRV::OpLifetimeStop;
5667 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5668 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5677 case Intrinsic::spv_saturate:
5678 return selectSaturate(ResVReg, ResType,
I);
5679 case Intrinsic::spv_nclamp:
5680 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5681 case Intrinsic::spv_uclamp:
5682 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5683 case Intrinsic::spv_sclamp:
5684 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5685 case Intrinsic::spv_subgroup_prefix_bit_count:
5686 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5687 case Intrinsic::spv_wave_active_countbits:
5688 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5689 case Intrinsic::spv_wave_all_equal:
5690 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5691 case Intrinsic::spv_wave_all:
5692 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5693 case Intrinsic::spv_wave_any:
5694 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5695 case Intrinsic::spv_subgroup_ballot:
5696 return selectWaveOpInst(ResVReg, ResType,
I,
5697 SPIRV::OpGroupNonUniformBallot);
5698 case Intrinsic::spv_wave_is_first_lane:
5699 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5700 case Intrinsic::spv_wave_reduce_or:
5701 return selectWaveReduceOp(ResVReg, ResType,
I,
5702 SPIRV::OpGroupNonUniformBitwiseOr);
5703 case Intrinsic::spv_wave_reduce_xor:
5704 return selectWaveReduceOp(ResVReg, ResType,
I,
5705 SPIRV::OpGroupNonUniformBitwiseXor);
5706 case Intrinsic::spv_wave_reduce_and:
5707 return selectWaveReduceOp(ResVReg, ResType,
I,
5708 SPIRV::OpGroupNonUniformBitwiseAnd);
5709 case Intrinsic::spv_wave_reduce_umax:
5710 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5711 case Intrinsic::spv_wave_reduce_max:
5712 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5713 case Intrinsic::spv_wave_reduce_umin:
5714 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5715 case Intrinsic::spv_wave_reduce_min:
5716 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5717 case Intrinsic::spv_wave_reduce_sum:
5718 return selectWaveReduceSum(ResVReg, ResType,
I);
5719 case Intrinsic::spv_wave_product:
5720 return selectWaveReduceProduct(ResVReg, ResType,
I);
5721 case Intrinsic::spv_wave_readlane:
5722 return selectWaveOpInst(ResVReg, ResType,
I,
5723 SPIRV::OpGroupNonUniformShuffle);
5724 case Intrinsic::spv_wave_prefix_sum:
5725 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5726 case Intrinsic::spv_wave_prefix_product:
5727 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5728 case Intrinsic::spv_quad_read_across_x: {
5729 return selectQuadSwap(ResVReg, ResType,
I, 0);
5731 case Intrinsic::spv_quad_read_across_y: {
5732 return selectQuadSwap(ResVReg, ResType,
I, 1);
5734 case Intrinsic::spv_quad_read_across_diagonal: {
5735 return selectQuadSwap(ResVReg, ResType,
I, 2);
5737 case Intrinsic::spv_radians:
5738 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5742 case Intrinsic::instrprof_increment:
5743 case Intrinsic::instrprof_increment_step:
5744 case Intrinsic::instrprof_value_profile:
5747 case Intrinsic::spv_value_md:
5749 case Intrinsic::spv_resource_handlefrombinding: {
5750 return selectHandleFromBinding(ResVReg, ResType,
I);
5752 case Intrinsic::spv_resource_counterhandlefrombinding:
5753 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5754 case Intrinsic::spv_resource_updatecounter:
5755 return selectUpdateCounter(ResVReg, ResType,
I);
5756 case Intrinsic::spv_resource_store_typedbuffer: {
5757 return selectImageWriteIntrinsic(
I);
5759 case Intrinsic::spv_resource_load_typedbuffer: {
5760 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5762 case Intrinsic::spv_resource_load_level: {
5763 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5765 case Intrinsic::spv_resource_getdimensions_x:
5766 case Intrinsic::spv_resource_getdimensions_xy:
5767 case Intrinsic::spv_resource_getdimensions_xyz: {
5768 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5770 case Intrinsic::spv_resource_getdimensions_levels_x:
5771 case Intrinsic::spv_resource_getdimensions_levels_xy:
5772 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5773 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5775 case Intrinsic::spv_resource_getdimensions_ms_xy:
5776 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5777 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5779 case Intrinsic::spv_resource_calculate_lod:
5780 case Intrinsic::spv_resource_calculate_lod_unclamped:
5781 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5782 case Intrinsic::spv_resource_sample:
5783 case Intrinsic::spv_resource_sample_clamp:
5784 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5785 case Intrinsic::spv_resource_samplebias:
5786 case Intrinsic::spv_resource_samplebias_clamp:
5787 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5788 case Intrinsic::spv_resource_samplegrad:
5789 case Intrinsic::spv_resource_samplegrad_clamp:
5790 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5791 case Intrinsic::spv_resource_samplelevel:
5792 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5793 case Intrinsic::spv_resource_samplecmp:
5794 case Intrinsic::spv_resource_samplecmp_clamp:
5795 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5796 case Intrinsic::spv_resource_samplecmplevelzero:
5797 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5798 case Intrinsic::spv_resource_gather:
5799 case Intrinsic::spv_resource_gather_cmp:
5800 return selectGatherIntrinsic(ResVReg, ResType,
I);
5801 case Intrinsic::spv_resource_getbasepointer:
5802 case Intrinsic::spv_resource_getpointer: {
5803 return selectResourceGetPointer(ResVReg, ResType,
I);
5805 case Intrinsic::spv_pushconstant_getpointer: {
5806 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5808 case Intrinsic::spv_discard: {
5809 return selectDiscard(ResVReg, ResType,
I);
5811 case Intrinsic::spv_resource_nonuniformindex: {
5812 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5814 case Intrinsic::spv_unpackhalf2x16: {
5815 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5817 case Intrinsic::spv_packhalf2x16: {
5818 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5820 case Intrinsic::spv_ddx:
5821 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5822 case Intrinsic::spv_ddy:
5823 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5824 case Intrinsic::spv_ddx_coarse:
5825 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5826 case Intrinsic::spv_ddy_coarse:
5827 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5828 case Intrinsic::spv_ddx_fine:
5829 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5830 case Intrinsic::spv_ddy_fine:
5831 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5832 case Intrinsic::spv_fwidth:
5833 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5834 case Intrinsic::spv_masked_gather:
5835 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5836 return selectMaskedGather(ResVReg, ResType,
I);
5837 return diagnoseUnsupported(
5838 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5839 case Intrinsic::spv_masked_scatter:
5840 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5841 return selectMaskedScatter(
I);
5842 return diagnoseUnsupported(
5843 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5844 case Intrinsic::returnaddress:
5845 case Intrinsic::frameaddress: {
5847 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5854 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5859bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5860 SPIRVTypeInst ResType,
5861 MachineInstr &
I)
const {
5864 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5871bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5872 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5874 assert(Intr.getIntrinsicID() ==
5875 Intrinsic::spv_resource_counterhandlefrombinding);
5878 Register MainHandleReg = Intr.getOperand(2).getReg();
5880 assert(MainHandleDef->getIntrinsicID() ==
5881 Intrinsic::spv_resource_handlefrombinding);
5885 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5886 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5887 std::string CounterName =
5892 MachineIRBuilder MIRBuilder(
I);
5894 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5896 ArraySize, IndexReg, CounterName, MIRBuilder);
5898 return BuildCOPY(ResVReg, CounterVarReg,
I);
5901bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5902 SPIRVTypeInst ResType,
5903 MachineInstr &
I)
const {
5905 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5907 Register CounterHandleReg = Intr.getOperand(2).getReg();
5908 Register IncrReg = Intr.getOperand(3).getReg();
5915 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5916 assert(CounterVarPointeeType &&
5917 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5918 "Counter variable must be a struct");
5920 SPIRV::StorageClass::StorageBuffer &&
5921 "Counter variable must be in the storage buffer storage class");
5923 "Counter variable must have exactly 1 member in the struct");
5924 const SPIRVTypeInst MemberType =
5927 "Counter variable struct must have a single i32 member");
5931 MachineIRBuilder MIRBuilder(
I);
5933 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5936 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5942 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5945 .
addUse(CounterHandleReg)
5952 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5955 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5958 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5967 return BuildCOPY(ResVReg, AtomicRes,
I);
5975 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5983bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5984 SPIRVTypeInst ResType,
5985 MachineInstr &
I)
const {
5993 Register ImageReg =
I.getOperand(2).getReg();
6001 Register IdxReg =
I.getOperand(3).getReg();
6003 MachineInstr &Pos =
I;
6005 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
6009bool SPIRVInstructionSelector::generateSampleImage(
6012 DebugLoc Loc, MachineInstr &Pos)
const {
6023 if (!loadHandleBeforePosition(NewSamplerReg,
6029 MachineIRBuilder MIRBuilder(Pos);
6042 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
6043 ImOps.Lod.has_value();
6044 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
6045 : SPIRV::OpImageSampleImplicitLod;
6047 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
6048 : SPIRV::OpImageSampleDrefImplicitLod;
6057 MIB.
addUse(*ImOps.Compare);
6059 uint32_t ImageOperands = 0;
6061 ImageOperands |= SPIRV::ImageOperand::Bias;
6063 ImageOperands |= SPIRV::ImageOperand::Lod;
6064 if (ImOps.GradX && ImOps.GradY)
6065 ImageOperands |= SPIRV::ImageOperand::Grad;
6066 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
6068 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6071 "Non-constant offsets are not supported in sample instructions.");
6076 ImageOperands |= SPIRV::ImageOperand::MinLod;
6078 if (ImageOperands != 0) {
6079 MIB.
addImm(ImageOperands);
6080 if (ImageOperands & SPIRV::ImageOperand::Bias)
6082 if (ImageOperands & SPIRV::ImageOperand::Lod)
6084 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6085 MIB.
addUse(*ImOps.GradX);
6086 MIB.
addUse(*ImOps.GradY);
6089 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6090 MIB.
addUse(*ImOps.Offset);
6091 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6092 MIB.
addUse(*ImOps.MinLod);
6099bool SPIRVInstructionSelector::selectImageQuerySize(
6101 std::optional<Register> LodReg)
const {
6103 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6106 "ImageReg is not an image type.");
6108 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6110 unsigned NumComponents = 0;
6112 case SPIRV::Dim::DIM_1D:
6113 case SPIRV::Dim::DIM_Buffer:
6114 NumComponents =
IsArray ? 2 : 1;
6116 case SPIRV::Dim::DIM_2D:
6117 case SPIRV::Dim::DIM_Cube:
6118 case SPIRV::Dim::DIM_Rect:
6119 NumComponents =
IsArray ? 3 : 2;
6121 case SPIRV::Dim::DIM_3D:
6125 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6130 SPIRVTypeInst ResType =
6135 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6145bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6146 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6147 Register ImageReg =
I.getOperand(2).getReg();
6154 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6157bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6158 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6159 Register ImageReg =
I.getOperand(2).getReg();
6168 Register LodReg =
I.getOperand(3).getReg();
6171 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6173 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6180 TII.get(SPIRV::OpImageQueryLevels))
6187 TII.get(SPIRV::OpCompositeConstruct))
6197bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6198 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6199 Register ImageReg =
I.getOperand(2).getReg();
6210 "OpImageQuerySamples requires a multisampled image");
6212 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6220 TII.get(SPIRV::OpImageQuerySamples))
6227 TII.get(SPIRV::OpCompositeConstruct))
6237bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6238 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6239 Register ImageReg =
I.getOperand(2).getReg();
6240 Register SamplerReg =
I.getOperand(3).getReg();
6241 Register CoordinateReg =
I.getOperand(4).getReg();
6257 if (!loadHandleBeforePosition(
6262 MachineIRBuilder MIRBuilder(
I);
6268 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6278 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6285 unsigned ExtractedIndex =
6287 Intrinsic::spv_resource_calculate_lod_unclamped
6291 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6292 TII.get(SPIRV::OpCompositeExtract))
6302bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6303 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6304 Register ImageReg =
I.getOperand(2).getReg();
6305 Register SamplerReg =
I.getOperand(3).getReg();
6306 Register CoordinateReg =
I.getOperand(4).getReg();
6307 ImageOperands ImOps;
6308 if (
I.getNumOperands() > 5)
6309 ImOps.Offset =
I.getOperand(5).getReg();
6310 if (
I.getNumOperands() > 6)
6311 ImOps.MinLod =
I.getOperand(6).getReg();
6312 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6313 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6316bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6317 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6318 Register ImageReg =
I.getOperand(2).getReg();
6319 Register SamplerReg =
I.getOperand(3).getReg();
6320 Register CoordinateReg =
I.getOperand(4).getReg();
6321 ImageOperands ImOps;
6322 ImOps.Bias =
I.getOperand(5).getReg();
6323 if (
I.getNumOperands() > 6)
6324 ImOps.Offset =
I.getOperand(6).getReg();
6325 if (
I.getNumOperands() > 7)
6326 ImOps.MinLod =
I.getOperand(7).getReg();
6327 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6328 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6331bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6332 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6333 Register ImageReg =
I.getOperand(2).getReg();
6334 Register SamplerReg =
I.getOperand(3).getReg();
6335 Register CoordinateReg =
I.getOperand(4).getReg();
6336 ImageOperands ImOps;
6337 ImOps.GradX =
I.getOperand(5).getReg();
6338 ImOps.GradY =
I.getOperand(6).getReg();
6339 if (
I.getNumOperands() > 7)
6340 ImOps.Offset =
I.getOperand(7).getReg();
6341 if (
I.getNumOperands() > 8)
6342 ImOps.MinLod =
I.getOperand(8).getReg();
6343 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6344 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6347bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6348 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6349 Register ImageReg =
I.getOperand(2).getReg();
6350 Register SamplerReg =
I.getOperand(3).getReg();
6351 Register CoordinateReg =
I.getOperand(4).getReg();
6352 ImageOperands ImOps;
6353 ImOps.Lod =
I.getOperand(5).getReg();
6354 if (
I.getNumOperands() > 6)
6355 ImOps.Offset =
I.getOperand(6).getReg();
6356 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6357 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6360bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6361 SPIRVTypeInst ResType,
6362 MachineInstr &
I)
const {
6363 Register ImageReg =
I.getOperand(2).getReg();
6364 Register SamplerReg =
I.getOperand(3).getReg();
6365 Register CoordinateReg =
I.getOperand(4).getReg();
6366 ImageOperands ImOps;
6367 ImOps.Compare =
I.getOperand(5).getReg();
6368 if (
I.getNumOperands() > 6)
6369 ImOps.Offset =
I.getOperand(6).getReg();
6370 if (
I.getNumOperands() > 7)
6371 ImOps.MinLod =
I.getOperand(7).getReg();
6372 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6373 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6376bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6377 SPIRVTypeInst ResType,
6378 MachineInstr &
I)
const {
6379 Register ImageReg =
I.getOperand(2).getReg();
6380 Register CoordinateReg =
I.getOperand(3).getReg();
6381 Register LodReg =
I.getOperand(4).getReg();
6383 ImageOperands ImOps;
6385 if (
I.getNumOperands() > 5)
6386 ImOps.Offset =
I.getOperand(5).getReg();
6398 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6399 I.getDebugLoc(),
I, &ImOps);
6402bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6403 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6404 Register ImageReg =
I.getOperand(2).getReg();
6405 Register SamplerReg =
I.getOperand(3).getReg();
6406 Register CoordinateReg =
I.getOperand(4).getReg();
6407 ImageOperands ImOps;
6408 ImOps.Compare =
I.getOperand(5).getReg();
6409 if (
I.getNumOperands() > 6)
6410 ImOps.Offset =
I.getOperand(6).getReg();
6413 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6414 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6417bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6418 SPIRVTypeInst ResType,
6419 MachineInstr &
I)
const {
6420 Register ImageReg =
I.getOperand(2).getReg();
6421 Register SamplerReg =
I.getOperand(3).getReg();
6422 Register CoordinateReg =
I.getOperand(4).getReg();
6425 "ImageReg is not an image type.");
6430 ComponentOrCompareReg =
I.getOperand(5).getReg();
6431 OffsetReg =
I.getOperand(6).getReg();
6434 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6438 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6439 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6440 Dim != SPIRV::Dim::DIM_Rect) {
6442 "Gather operations are only supported for 2D, Cube, and Rect images.");
6449 if (!loadHandleBeforePosition(
6454 MachineIRBuilder MIRBuilder(
I);
6455 SPIRVTypeInst SampledImageType =
6460 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6468 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6470 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6472 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6477 .
addUse(ComponentOrCompareReg);
6479 uint32_t ImageOperands = 0;
6480 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6481 if (Dim == SPIRV::Dim::DIM_Cube) {
6483 "Gather operations with offset are not supported for Cube images.");
6487 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6489 ImageOperands |= SPIRV::ImageOperand::Offset;
6493 if (ImageOperands != 0) {
6494 MIB.
addImm(ImageOperands);
6496 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6504bool SPIRVInstructionSelector::generateImageReadOrFetch(
6507 const ImageOperands *ImOps)
const {
6510 "ImageReg is not an image type.");
6512 bool IsSignedInteger =
6517 bool IsFetch = (SampledOp.getImm() == 1);
6519 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6520 uint32_t ImageOperandsMask = 0;
6521 if (IsSignedInteger)
6522 ImageOperandsMask |= 0x1000;
6524 if (IsFetch && ImOps) {
6526 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6527 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6529 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6531 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6535 if (ImageOperandsMask != 0) {
6536 MIB.
addImm(ImageOperandsMask);
6537 if (IsFetch && ImOps) {
6540 if (ImOps->Offset &&
6541 (ImageOperandsMask &
6542 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6543 MIB.
addUse(*ImOps->Offset);
6552 SPIRVTypeInst SampledType =
6555 SPIRVTypeInst ReadType =
6556 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6557 bool ReadTypeMatchesResult = ReadType == ResType;
6559 Register ReadReg = ReadTypeMatchesResult
6565 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6571 BMI.constrainAllUses(
TII,
TRI, RBI);
6573 if (ReadTypeMatchesResult)
6586 if (ResultSize == 1) {
6595 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6598bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6599 SPIRVTypeInst ResType,
6600 MachineInstr &
I)
const {
6601 Register ResourcePtr =
I.getOperand(2).getReg();
6603 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6612 MachineIRBuilder MIRBuilder(
I);
6617 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6623 if (
I.getNumExplicitOperands() > 3) {
6624 Register IndexReg =
I.getOperand(3).getReg();
6631bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6632 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6637bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6638 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6639 Register ObjReg =
I.getOperand(2).getReg();
6640 if (!BuildCOPY(ResVReg, ObjReg,
I))
6650 decorateUsesAsNonUniform(ResVReg);
6654void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6657 {NonUniformReg,
nullptr}};
6658 llvm::SmallSet<Register, 8> Visited;
6659 while (WorkList.
size() > 0) {
6662 if (!Visited.
insert(CurrentReg).second)
6665 bool IsDecorated =
false;
6667 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6668 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6674 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6676 if (ResultReg == CurrentReg)
6684 MachineInstr &InsertPt =
6687 SPIRV::Decoration::NonUniformEXT, {});
6692bool SPIRVInstructionSelector::extractSubvector(
6694 MachineInstr &InsertionPoint)
const {
6696 [[maybe_unused]]
uint64_t InputSize =
6699 [[maybe_unused]]
bool IsLongVectorEXT =
6701 assert((InputSize > 1 || IsLongVectorEXT) &&
"The input must be a vector.");
6702 assert((ResultSize > 1 || IsLongVectorEXT) &&
"The result must be a vector.");
6703 assert(ResultSize < InputSize &&
6704 "Cannot extract more element than there are in the input.");
6711 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6720 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6722 TII.get(SPIRV::OpCompositeConstruct))
6726 for (
Register ComponentReg : ComponentRegisters)
6727 MIB.
addUse(ComponentReg);
6732bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6733 MachineInstr &
I)
const {
6740 Register ImageReg =
I.getOperand(1).getReg();
6748 Register CoordinateReg =
I.getOperand(2).getReg();
6749 Register DataReg =
I.getOperand(3).getReg();
6752 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6760Register SPIRVInstructionSelector::buildPointerToResource(
6761 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6762 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6763 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6765 if (ArraySize == 1) {
6766 SPIRVTypeInst PtrType =
6769 "SpirvResType did not have an explicit layout.");
6774 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6775 SPIRVTypeInst VarPointerType =
6778 VarPointerType, Set,
Binding, Name, MIRBuilder);
6780 SPIRVTypeInst ResPointerType =
6793bool SPIRVInstructionSelector::selectFirstBitSet16(
6794 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6795 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6797 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6801 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6804bool SPIRVInstructionSelector::selectFirstBitSet32(
6806 unsigned BitSetOpcode)
const {
6807 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6810 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6817bool SPIRVInstructionSelector::selectFirstBitSet64(
6819 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6833 if (ComponentCount > 2) {
6834 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6836 unsigned Opcode) ->
bool {
6837 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6841 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6845 MachineIRBuilder MIRBuilder(
I);
6847 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6851 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6857 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6867 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6868 SPIRV::OpVectorExtractDynamic))
6870 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6871 SPIRV::OpVectorExtractDynamic))
6875 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6876 TII.get(SPIRV::OpVectorShuffle))
6884 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6890 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6891 TII.get(SPIRV::OpVectorShuffle))
6899 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6919 SelectOp = SPIRV::OpSelectSISCond;
6920 AddOp = SPIRV::OpIAddS;
6928 SelectOp = SPIRV::OpSelectVIVCond;
6929 AddOp = SPIRV::OpIAddV;
6935 Register RegSecondaryOffset = Reg0;
6939 if (SwapPrimarySide) {
6940 PrimaryReg = LowReg;
6941 SecondaryReg = HighReg;
6942 RegPrimaryOffset = Reg0;
6943 RegSecondaryOffset = Reg32;
6948 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6949 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6954 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6955 SPIRV::OpINotEqual))
6962 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6963 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6968 if (SwapPrimarySide) {
6970 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6971 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6982 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6983 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6988 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6989 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6992 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6996bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6997 SPIRVTypeInst ResType,
6999 bool IsSigned)
const {
7001 Register OpReg =
I.getOperand(2).getReg();
7004 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
7005 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
7009 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7011 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7013 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7016 return diagnoseUnsupported(
7018 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
7022bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
7023 SPIRVTypeInst ResType,
7024 MachineInstr &
I)
const {
7026 Register OpReg =
I.getOperand(2).getReg();
7031 unsigned ExtendOpcode = SPIRV::OpUConvert;
7032 unsigned BitSetOpcode = GL::FindILsb;
7036 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7038 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7040 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7043 return diagnoseUnsupported(
I,
7044 "spv_firstbitlow only supports 16,32,64 bits.");
7048bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
7049 SPIRVTypeInst ResType,
7050 MachineInstr &
I)
const {
7054 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
7057 .
addUse(
I.getOperand(2).getReg())
7060 unsigned Alignment =
I.getOperand(3).getImm();
7074 while (!Worklist.
empty()) {
7076 switch (
T->getOpcode()) {
7077 case SPIRV::OpTypeInt:
7078 case SPIRV::OpTypeFloat:
7079 case SPIRV::OpTypePointer:
7081 case SPIRV::OpTypeVector:
7082 case SPIRV::OpTypeVectorIdEXT:
7083 case SPIRV::OpTypeMatrix:
7084 case SPIRV::OpTypeArray: {
7085 Register OperandReg =
T->getOperand(1).getReg();
7089 case SPIRV::OpTypeStruct:
7090 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7091 Register OperandReg =
T->getOperand(Idx).getReg();
7106 Register TypeReg = Ty->getOperand(0).getReg();
7107 if (!Visited.
insert(TypeReg).second)
7110 switch (Ty->getOpcode()) {
7111 case SPIRV::OpTypePointer:
7112 if (Ty->getOperand(1).getImm() == SPIRV::StorageClass::StorageBuffer)
7116 case SPIRV::OpTypeArray:
7117 case SPIRV::OpTypeRuntimeArray:
7120 case SPIRV::OpTypeStruct:
7121 for (
unsigned I = 1;
I < Ty->getNumOperands(); ++
I)
7137bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7138 assert(
I.getNumExplicitOperands() == 2);
7140 Register MsgReg =
I.getOperand(1).getReg();
7142 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7145 return diagnoseUnsupported(
7147 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7148 "scalar, pointer, vector, matrix, or aggregate of such types)");
7151 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7158bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7167 uint32_t MsgVal = ~0
u;
7168 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7169 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7172 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7175 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7182bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7183 SPIRVTypeInst ResType,
7184 MachineInstr &
I)
const {
7191 bool UseUntypedPointers =
7192 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7194 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7197 MachineIRBuilder MIRBuilder(
I);
7200 .
addImm(SPIRV::Extension::SPV_KHR_variable_pointers);
7202 .
addImm(SPIRV::Capability::VariablePointersStorageBuffer);
7205 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7208 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7212 if (UseUntypedPointers) {
7216 return diagnoseUnsupported(
7217 I,
"could not deduce the data type of an untyped variable");
7223 unsigned Alignment =
I.getOperand(2).getImm();
7230bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7235 const MachineInstr *PrevI =
I.getPrevNode();
7237 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7241 .
addMBB(
I.getOperand(0).getMBB())
7246 .
addMBB(
I.getOperand(0).getMBB())
7251bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7262 const MachineInstr *NextI =
I.getNextNode();
7264 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7270 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7272 .
addUse(
I.getOperand(0).getReg())
7273 .
addMBB(
I.getOperand(1).getMBB())
7279bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7280 MachineInstr &
I)
const {
7282 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7284 const unsigned NumOps =
I.getNumOperands();
7285 for (
unsigned i = 1; i <
NumOps; i += 2) {
7286 MIB.
addUse(
I.getOperand(i + 0).getReg());
7287 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7293bool SPIRVInstructionSelector::selectGlobalValue(
7294 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7296 MachineIRBuilder MIRBuilder(
I);
7297 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7300 std::string GlobalIdent;
7302 unsigned &
ID = UnnamedGlobalIDs[GV];
7304 ID = UnnamedGlobalIDs.
size();
7305 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7331 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7338 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7343 MachineInstrBuilder MIB1 =
7344 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7347 MachineInstrBuilder MIB2 =
7349 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7353 GR.
add(ConstVal, MIB2);
7361 MachineInstrBuilder MIB3 =
7362 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7365 GR.
add(ConstVal, MIB3);
7371 assert(NewReg != ResVReg);
7372 return BuildCOPY(ResVReg, NewReg,
I);
7382 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7385 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7391 SPIRVTypeInst ResType =
7395 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7400 if (
GlobalVar->isExternallyInitialized() &&
7401 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7402 constexpr unsigned ReadWriteINTEL = 3u;
7405 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7411bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7412 SPIRVTypeInst ResType,
7413 MachineInstr &
I)
const {
7415 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7423 MachineIRBuilder MIRBuilder(
I);
7428 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7431 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7433 .
add(
I.getOperand(1))
7447 APFloat::rmNearestTiesToEven, &LosesInfo);
7452 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
7462bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7463 SPIRVTypeInst ResType,
7464 MachineInstr &
I)
const {
7467 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7473 Register ExpReg =
I.getOperand(2).getReg();
7475 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7476 SPIRV::OpConvertSToF))
7478 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7485bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7486 SPIRVTypeInst ResType,
7487 MachineInstr &
I)
const {
7503 MachineIRBuilder MIRBuilder(
I);
7504 SPIRVTypeInst FloatType =
7508 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7521 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7522 const bool IsUntyped =
7523 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7525 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7526 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7527 : SPIRV::OpVariable))
7530 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7538 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7541 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7544 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7548 Register IntegralPartReg =
I.getOperand(1).getReg();
7551 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7561 assert(
false &&
"GLSL::Modf is deprecated.");
7572bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7573 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7574 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7575 MachineIRBuilder MIRBuilder(
I);
7576 const SPIRVTypeInst Vec3Ty =
7579 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7591 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7595 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7601 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7608 assert(
I.getOperand(2).isReg());
7609 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7613 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7624bool SPIRVInstructionSelector::loadBuiltinInputID(
7625 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7626 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7627 MachineIRBuilder MIRBuilder(
I);
7629 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7644 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7648 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7657SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7658 MachineInstr &
I)
const {
7659 MachineIRBuilder MIRBuilder(
I);
7670bool SPIRVInstructionSelector::loadHandleBeforePosition(
7671 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7672 MachineInstr &Pos)
const {
7675 Intrinsic::spv_resource_handlefrombinding);
7683 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7684 MachineIRBuilder MIRBuilder(HandleDef);
7685 SPIRVTypeInst VarType = ResType;
7686 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7688 if (IsStructuredBuffer) {
7697 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7700 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7701 ArraySize, IndexReg, Name, MIRBuilder);
7705 uint32_t LoadOpcode =
7706 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7716bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7717 MachineInstr &
I)
const {
7719 return diagnoseUnsupported(
7720 I,
"this instruction is only supported in shaders.");
7725InstructionSelector *
7729 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
MachineInstrBuilder & UseMI
MachineInstrBuilder MachineInstrBuilder & DefMI
#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
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 containsStorageBufferPointer(SPIRVTypeInst Ty, const SPIRVGlobalRegistry &GR, SmallSet< Register, 8 > &Visited)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef)
static unsigned getPtrCmpOpcode(unsigned Pred)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
static Register convertPtrToInt(Register Reg, LLT ConvTy, SPIRVTypeInst SpvType, LegalizerHelper &Helper, MachineRegisterInfo &MRI, SPIRVGlobalRegistry *GR)
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
This file defines the SmallSet class.
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.
static APInt getSignMask(unsigned BitWidth)
Get the SignMask for a specific bit width.
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.
static LLVM_ABI std::optional< GFConstant > getConstant(Register Const, const MachineRegisterInfo &MRI)
static LLVM_ABI std::optional< GIConstant > getConstant(Register Const, const MachineRegisterInfo &MRI)
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.
instr_iterator instr_end()
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 substituteRegister(Register FromReg, Register ToReg, unsigned SubIdx, const TargetRegisterInfo &RegInfo)
Replace all occurrences of FromReg with ToReg:SubIdx, properly composing subreg indices where necessa...
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 LLVM_READONLY 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.
bool hasOneUse(Register RegNo) const
hasOneUse - Return true if there is exactly one instruction using the specified 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)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC, bool ForceTyped=false)
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
SPIRVTypeInst getOpTypeVoid(MachineIRBuilder &MIRBuilder)
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
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
SPIRVTypeInst getUntypedPtrElementType(Register Reg) 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 isAnyTypeFloat() 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.
SmallSet - This maintains a set of unique values, optimizing for the case when the set is small (less...
std::pair< const_iterator, bool > insert(const T &V)
insert - Insert an element into the set if it isn't already there.
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.
An efficient, type-erasing, non-owning reference to a callable.
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.
constexpr uint64_t PointerSize
aarch64 pointer size.
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.
unsigned getOpcode(const VPValue *V)
Return the instruction opcode for the recipe defining V or 0 for unsupported recipes and VPValues not...
This is an optimization pass for GlobalISel generic memory operations.
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)
uint32_t getMemSemanticsWithStorageClass(const Triple &TT, uint32_t OrderSem, uint32_t StorageClassSem)
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.
iterator_range< T > make_range(T x, T y)
Convenience function for iterating over sub-ranges.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
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...
SPIRV::Scope::Scope getMemScope(const Triple &TT, LLVMContext &Ctx, SyncScope::ID Id)
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)
bool isVectorType(SPIRVTypeInst SPVTy)
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.
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