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_readlane_first:
5725 return selectWaveOpInst(ResVReg, ResType,
I,
5726 SPIRV::OpGroupNonUniformBroadcastFirst);
5727 case Intrinsic::spv_wave_prefix_sum:
5728 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5729 case Intrinsic::spv_wave_prefix_product:
5730 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5731 case Intrinsic::spv_quad_read_across_x: {
5732 return selectQuadSwap(ResVReg, ResType,
I, 0);
5734 case Intrinsic::spv_quad_read_across_y: {
5735 return selectQuadSwap(ResVReg, ResType,
I, 1);
5737 case Intrinsic::spv_quad_read_across_diagonal: {
5738 return selectQuadSwap(ResVReg, ResType,
I, 2);
5740 case Intrinsic::spv_radians:
5741 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5745 case Intrinsic::instrprof_increment:
5746 case Intrinsic::instrprof_increment_step:
5747 case Intrinsic::instrprof_value_profile:
5750 case Intrinsic::spv_value_md:
5752 case Intrinsic::spv_resource_handlefrombinding: {
5753 return selectHandleFromBinding(ResVReg, ResType,
I);
5755 case Intrinsic::spv_resource_counterhandlefrombinding:
5756 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5757 case Intrinsic::spv_resource_updatecounter:
5758 return selectUpdateCounter(ResVReg, ResType,
I);
5759 case Intrinsic::spv_resource_store_typedbuffer: {
5760 return selectImageWriteIntrinsic(
I);
5762 case Intrinsic::spv_resource_load_typedbuffer: {
5763 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5765 case Intrinsic::spv_resource_load_level: {
5766 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5768 case Intrinsic::spv_resource_getdimensions_x:
5769 case Intrinsic::spv_resource_getdimensions_xy:
5770 case Intrinsic::spv_resource_getdimensions_xyz: {
5771 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5773 case Intrinsic::spv_resource_getdimensions_levels_x:
5774 case Intrinsic::spv_resource_getdimensions_levels_xy:
5775 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5776 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5778 case Intrinsic::spv_resource_getdimensions_ms_xy:
5779 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5780 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5782 case Intrinsic::spv_resource_calculate_lod:
5783 case Intrinsic::spv_resource_calculate_lod_unclamped:
5784 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5785 case Intrinsic::spv_resource_sample:
5786 case Intrinsic::spv_resource_sample_clamp:
5787 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5788 case Intrinsic::spv_resource_samplebias:
5789 case Intrinsic::spv_resource_samplebias_clamp:
5790 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5791 case Intrinsic::spv_resource_samplegrad:
5792 case Intrinsic::spv_resource_samplegrad_clamp:
5793 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5794 case Intrinsic::spv_resource_samplelevel:
5795 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5796 case Intrinsic::spv_resource_samplecmp:
5797 case Intrinsic::spv_resource_samplecmp_clamp:
5798 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5799 case Intrinsic::spv_resource_samplecmplevelzero:
5800 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5801 case Intrinsic::spv_resource_gather:
5802 case Intrinsic::spv_resource_gather_cmp:
5803 return selectGatherIntrinsic(ResVReg, ResType,
I);
5804 case Intrinsic::spv_resource_getbasepointer:
5805 case Intrinsic::spv_resource_getpointer: {
5806 return selectResourceGetPointer(ResVReg, ResType,
I);
5808 case Intrinsic::spv_pushconstant_getpointer: {
5809 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5811 case Intrinsic::spv_discard: {
5812 return selectDiscard(ResVReg, ResType,
I);
5814 case Intrinsic::spv_resource_nonuniformindex: {
5815 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5817 case Intrinsic::spv_unpackhalf2x16: {
5818 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5820 case Intrinsic::spv_packhalf2x16: {
5821 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5823 case Intrinsic::spv_ddx:
5824 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5825 case Intrinsic::spv_ddy:
5826 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5827 case Intrinsic::spv_ddx_coarse:
5828 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5829 case Intrinsic::spv_ddy_coarse:
5830 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5831 case Intrinsic::spv_ddx_fine:
5832 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5833 case Intrinsic::spv_ddy_fine:
5834 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5835 case Intrinsic::spv_fwidth:
5836 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5837 case Intrinsic::spv_masked_gather:
5838 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5839 return selectMaskedGather(ResVReg, ResType,
I);
5840 return diagnoseUnsupported(
5841 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5842 case Intrinsic::spv_masked_scatter:
5843 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5844 return selectMaskedScatter(
I);
5845 return diagnoseUnsupported(
5846 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5847 case Intrinsic::returnaddress:
5848 case Intrinsic::frameaddress: {
5850 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5857 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5862bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5863 SPIRVTypeInst ResType,
5864 MachineInstr &
I)
const {
5867 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5874bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5875 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5877 assert(Intr.getIntrinsicID() ==
5878 Intrinsic::spv_resource_counterhandlefrombinding);
5881 Register MainHandleReg = Intr.getOperand(2).getReg();
5883 assert(MainHandleDef->getIntrinsicID() ==
5884 Intrinsic::spv_resource_handlefrombinding);
5888 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5889 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5890 std::string CounterName =
5895 MachineIRBuilder MIRBuilder(
I);
5897 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5899 ArraySize, IndexReg, CounterName, MIRBuilder);
5901 return BuildCOPY(ResVReg, CounterVarReg,
I);
5904bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5905 SPIRVTypeInst ResType,
5906 MachineInstr &
I)
const {
5908 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5910 Register CounterHandleReg = Intr.getOperand(2).getReg();
5911 Register IncrReg = Intr.getOperand(3).getReg();
5918 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5919 assert(CounterVarPointeeType &&
5920 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5921 "Counter variable must be a struct");
5923 SPIRV::StorageClass::StorageBuffer &&
5924 "Counter variable must be in the storage buffer storage class");
5926 "Counter variable must have exactly 1 member in the struct");
5927 const SPIRVTypeInst MemberType =
5930 "Counter variable struct must have a single i32 member");
5934 MachineIRBuilder MIRBuilder(
I);
5936 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5939 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5945 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5948 .
addUse(CounterHandleReg)
5955 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5958 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5961 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5970 return BuildCOPY(ResVReg, AtomicRes,
I);
5978 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5986bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5987 SPIRVTypeInst ResType,
5988 MachineInstr &
I)
const {
5996 Register ImageReg =
I.getOperand(2).getReg();
6004 Register IdxReg =
I.getOperand(3).getReg();
6006 MachineInstr &Pos =
I;
6008 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
6012bool SPIRVInstructionSelector::generateSampleImage(
6015 DebugLoc Loc, MachineInstr &Pos)
const {
6026 if (!loadHandleBeforePosition(NewSamplerReg,
6032 MachineIRBuilder MIRBuilder(Pos);
6045 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
6046 ImOps.Lod.has_value();
6047 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
6048 : SPIRV::OpImageSampleImplicitLod;
6050 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
6051 : SPIRV::OpImageSampleDrefImplicitLod;
6060 MIB.
addUse(*ImOps.Compare);
6062 uint32_t ImageOperands = 0;
6064 ImageOperands |= SPIRV::ImageOperand::Bias;
6066 ImageOperands |= SPIRV::ImageOperand::Lod;
6067 if (ImOps.GradX && ImOps.GradY)
6068 ImageOperands |= SPIRV::ImageOperand::Grad;
6069 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
6071 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6074 "Non-constant offsets are not supported in sample instructions.");
6079 ImageOperands |= SPIRV::ImageOperand::MinLod;
6081 if (ImageOperands != 0) {
6082 MIB.
addImm(ImageOperands);
6083 if (ImageOperands & SPIRV::ImageOperand::Bias)
6085 if (ImageOperands & SPIRV::ImageOperand::Lod)
6087 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6088 MIB.
addUse(*ImOps.GradX);
6089 MIB.
addUse(*ImOps.GradY);
6092 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6093 MIB.
addUse(*ImOps.Offset);
6094 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6095 MIB.
addUse(*ImOps.MinLod);
6102bool SPIRVInstructionSelector::selectImageQuerySize(
6104 std::optional<Register> LodReg)
const {
6106 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6109 "ImageReg is not an image type.");
6111 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6113 unsigned NumComponents = 0;
6115 case SPIRV::Dim::DIM_1D:
6116 case SPIRV::Dim::DIM_Buffer:
6117 NumComponents =
IsArray ? 2 : 1;
6119 case SPIRV::Dim::DIM_2D:
6120 case SPIRV::Dim::DIM_Cube:
6121 case SPIRV::Dim::DIM_Rect:
6122 NumComponents =
IsArray ? 3 : 2;
6124 case SPIRV::Dim::DIM_3D:
6128 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6133 SPIRVTypeInst ResType =
6138 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6148bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6149 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6150 Register ImageReg =
I.getOperand(2).getReg();
6157 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6160bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6161 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6162 Register ImageReg =
I.getOperand(2).getReg();
6171 Register LodReg =
I.getOperand(3).getReg();
6174 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6176 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6183 TII.get(SPIRV::OpImageQueryLevels))
6190 TII.get(SPIRV::OpCompositeConstruct))
6200bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6201 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6202 Register ImageReg =
I.getOperand(2).getReg();
6213 "OpImageQuerySamples requires a multisampled image");
6215 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6223 TII.get(SPIRV::OpImageQuerySamples))
6230 TII.get(SPIRV::OpCompositeConstruct))
6240bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6241 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6242 Register ImageReg =
I.getOperand(2).getReg();
6243 Register SamplerReg =
I.getOperand(3).getReg();
6244 Register CoordinateReg =
I.getOperand(4).getReg();
6260 if (!loadHandleBeforePosition(
6265 MachineIRBuilder MIRBuilder(
I);
6271 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6281 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6288 unsigned ExtractedIndex =
6290 Intrinsic::spv_resource_calculate_lod_unclamped
6294 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6295 TII.get(SPIRV::OpCompositeExtract))
6305bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6306 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6307 Register ImageReg =
I.getOperand(2).getReg();
6308 Register SamplerReg =
I.getOperand(3).getReg();
6309 Register CoordinateReg =
I.getOperand(4).getReg();
6310 ImageOperands ImOps;
6311 if (
I.getNumOperands() > 5)
6312 ImOps.Offset =
I.getOperand(5).getReg();
6313 if (
I.getNumOperands() > 6)
6314 ImOps.MinLod =
I.getOperand(6).getReg();
6315 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6316 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6319bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6320 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6321 Register ImageReg =
I.getOperand(2).getReg();
6322 Register SamplerReg =
I.getOperand(3).getReg();
6323 Register CoordinateReg =
I.getOperand(4).getReg();
6324 ImageOperands ImOps;
6325 ImOps.Bias =
I.getOperand(5).getReg();
6326 if (
I.getNumOperands() > 6)
6327 ImOps.Offset =
I.getOperand(6).getReg();
6328 if (
I.getNumOperands() > 7)
6329 ImOps.MinLod =
I.getOperand(7).getReg();
6330 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6331 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6334bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6335 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6336 Register ImageReg =
I.getOperand(2).getReg();
6337 Register SamplerReg =
I.getOperand(3).getReg();
6338 Register CoordinateReg =
I.getOperand(4).getReg();
6339 ImageOperands ImOps;
6340 ImOps.GradX =
I.getOperand(5).getReg();
6341 ImOps.GradY =
I.getOperand(6).getReg();
6342 if (
I.getNumOperands() > 7)
6343 ImOps.Offset =
I.getOperand(7).getReg();
6344 if (
I.getNumOperands() > 8)
6345 ImOps.MinLod =
I.getOperand(8).getReg();
6346 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6347 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6350bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6351 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6352 Register ImageReg =
I.getOperand(2).getReg();
6353 Register SamplerReg =
I.getOperand(3).getReg();
6354 Register CoordinateReg =
I.getOperand(4).getReg();
6355 ImageOperands ImOps;
6356 ImOps.Lod =
I.getOperand(5).getReg();
6357 if (
I.getNumOperands() > 6)
6358 ImOps.Offset =
I.getOperand(6).getReg();
6359 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6360 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6363bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6364 SPIRVTypeInst ResType,
6365 MachineInstr &
I)
const {
6366 Register ImageReg =
I.getOperand(2).getReg();
6367 Register SamplerReg =
I.getOperand(3).getReg();
6368 Register CoordinateReg =
I.getOperand(4).getReg();
6369 ImageOperands ImOps;
6370 ImOps.Compare =
I.getOperand(5).getReg();
6371 if (
I.getNumOperands() > 6)
6372 ImOps.Offset =
I.getOperand(6).getReg();
6373 if (
I.getNumOperands() > 7)
6374 ImOps.MinLod =
I.getOperand(7).getReg();
6375 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6376 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6379bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6380 SPIRVTypeInst ResType,
6381 MachineInstr &
I)
const {
6382 Register ImageReg =
I.getOperand(2).getReg();
6383 Register CoordinateReg =
I.getOperand(3).getReg();
6384 Register LodReg =
I.getOperand(4).getReg();
6386 ImageOperands ImOps;
6388 if (
I.getNumOperands() > 5)
6389 ImOps.Offset =
I.getOperand(5).getReg();
6401 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6402 I.getDebugLoc(),
I, &ImOps);
6405bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6406 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6407 Register ImageReg =
I.getOperand(2).getReg();
6408 Register SamplerReg =
I.getOperand(3).getReg();
6409 Register CoordinateReg =
I.getOperand(4).getReg();
6410 ImageOperands ImOps;
6411 ImOps.Compare =
I.getOperand(5).getReg();
6412 if (
I.getNumOperands() > 6)
6413 ImOps.Offset =
I.getOperand(6).getReg();
6416 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6417 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6420bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6421 SPIRVTypeInst ResType,
6422 MachineInstr &
I)
const {
6423 Register ImageReg =
I.getOperand(2).getReg();
6424 Register SamplerReg =
I.getOperand(3).getReg();
6425 Register CoordinateReg =
I.getOperand(4).getReg();
6428 "ImageReg is not an image type.");
6433 ComponentOrCompareReg =
I.getOperand(5).getReg();
6434 OffsetReg =
I.getOperand(6).getReg();
6437 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6441 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6442 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6443 Dim != SPIRV::Dim::DIM_Rect) {
6445 "Gather operations are only supported for 2D, Cube, and Rect images.");
6452 if (!loadHandleBeforePosition(
6457 MachineIRBuilder MIRBuilder(
I);
6458 SPIRVTypeInst SampledImageType =
6463 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6471 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6473 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6475 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6480 .
addUse(ComponentOrCompareReg);
6482 uint32_t ImageOperands = 0;
6483 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6484 if (Dim == SPIRV::Dim::DIM_Cube) {
6486 "Gather operations with offset are not supported for Cube images.");
6490 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6492 ImageOperands |= SPIRV::ImageOperand::Offset;
6496 if (ImageOperands != 0) {
6497 MIB.
addImm(ImageOperands);
6499 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6507bool SPIRVInstructionSelector::generateImageReadOrFetch(
6510 const ImageOperands *ImOps)
const {
6513 "ImageReg is not an image type.");
6515 bool IsSignedInteger =
6520 bool IsFetch = (SampledOp.getImm() == 1);
6522 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6523 uint32_t ImageOperandsMask = 0;
6524 if (IsSignedInteger)
6525 ImageOperandsMask |= 0x1000;
6527 if (IsFetch && ImOps) {
6529 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6530 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6532 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6534 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6538 if (ImageOperandsMask != 0) {
6539 MIB.
addImm(ImageOperandsMask);
6540 if (IsFetch && ImOps) {
6543 if (ImOps->Offset &&
6544 (ImageOperandsMask &
6545 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6546 MIB.
addUse(*ImOps->Offset);
6555 SPIRVTypeInst SampledType =
6558 SPIRVTypeInst ReadType =
6559 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6560 bool ReadTypeMatchesResult = ReadType == ResType;
6562 Register ReadReg = ReadTypeMatchesResult
6568 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6574 BMI.constrainAllUses(
TII,
TRI, RBI);
6576 if (ReadTypeMatchesResult)
6589 if (ResultSize == 1) {
6598 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6601bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6602 SPIRVTypeInst ResType,
6603 MachineInstr &
I)
const {
6604 Register ResourcePtr =
I.getOperand(2).getReg();
6606 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6615 MachineIRBuilder MIRBuilder(
I);
6620 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6626 if (
I.getNumExplicitOperands() > 3) {
6627 Register IndexReg =
I.getOperand(3).getReg();
6634bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6635 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6640bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6641 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6642 Register ObjReg =
I.getOperand(2).getReg();
6643 if (!BuildCOPY(ResVReg, ObjReg,
I))
6653 decorateUsesAsNonUniform(ResVReg);
6657void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6660 {NonUniformReg,
nullptr}};
6661 llvm::SmallSet<Register, 8> Visited;
6662 while (WorkList.
size() > 0) {
6665 if (!Visited.
insert(CurrentReg).second)
6668 bool IsDecorated =
false;
6670 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6671 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6677 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6679 if (ResultReg == CurrentReg)
6687 MachineInstr &InsertPt =
6690 SPIRV::Decoration::NonUniformEXT, {});
6695bool SPIRVInstructionSelector::extractSubvector(
6697 MachineInstr &InsertionPoint)
const {
6699 [[maybe_unused]]
uint64_t InputSize =
6702 [[maybe_unused]]
bool IsLongVectorEXT =
6704 assert((InputSize > 1 || IsLongVectorEXT) &&
"The input must be a vector.");
6705 assert((ResultSize > 1 || IsLongVectorEXT) &&
"The result must be a vector.");
6706 assert(ResultSize < InputSize &&
6707 "Cannot extract more element than there are in the input.");
6714 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6723 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6725 TII.get(SPIRV::OpCompositeConstruct))
6729 for (
Register ComponentReg : ComponentRegisters)
6730 MIB.
addUse(ComponentReg);
6735bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6736 MachineInstr &
I)
const {
6743 Register ImageReg =
I.getOperand(1).getReg();
6751 Register CoordinateReg =
I.getOperand(2).getReg();
6752 Register DataReg =
I.getOperand(3).getReg();
6755 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6763Register SPIRVInstructionSelector::buildPointerToResource(
6764 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6765 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6766 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6768 if (ArraySize == 1) {
6769 SPIRVTypeInst PtrType =
6772 "SpirvResType did not have an explicit layout.");
6777 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6778 SPIRVTypeInst VarPointerType =
6781 VarPointerType, Set,
Binding, Name, MIRBuilder);
6783 SPIRVTypeInst ResPointerType =
6796bool SPIRVInstructionSelector::selectFirstBitSet16(
6797 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6798 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6800 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6804 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6807bool SPIRVInstructionSelector::selectFirstBitSet32(
6809 unsigned BitSetOpcode)
const {
6810 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6813 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6820bool SPIRVInstructionSelector::selectFirstBitSet64(
6822 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6836 if (ComponentCount > 2) {
6837 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6839 unsigned Opcode) ->
bool {
6840 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6844 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6848 MachineIRBuilder MIRBuilder(
I);
6850 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6854 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6860 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6870 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6871 SPIRV::OpVectorExtractDynamic))
6873 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6874 SPIRV::OpVectorExtractDynamic))
6878 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6879 TII.get(SPIRV::OpVectorShuffle))
6887 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6893 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6894 TII.get(SPIRV::OpVectorShuffle))
6902 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6922 SelectOp = SPIRV::OpSelectSISCond;
6923 AddOp = SPIRV::OpIAddS;
6931 SelectOp = SPIRV::OpSelectVIVCond;
6932 AddOp = SPIRV::OpIAddV;
6938 Register RegSecondaryOffset = Reg0;
6942 if (SwapPrimarySide) {
6943 PrimaryReg = LowReg;
6944 SecondaryReg = HighReg;
6945 RegPrimaryOffset = Reg0;
6946 RegSecondaryOffset = Reg32;
6951 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6952 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6957 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6958 SPIRV::OpINotEqual))
6965 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6966 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6971 if (SwapPrimarySide) {
6973 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6974 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6985 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6986 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6991 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6992 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6995 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6999bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
7000 SPIRVTypeInst ResType,
7002 bool IsSigned)
const {
7004 Register OpReg =
I.getOperand(2).getReg();
7007 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
7008 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
7012 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7014 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7016 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7019 return diagnoseUnsupported(
7021 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
7025bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
7026 SPIRVTypeInst ResType,
7027 MachineInstr &
I)
const {
7029 Register OpReg =
I.getOperand(2).getReg();
7034 unsigned ExtendOpcode = SPIRV::OpUConvert;
7035 unsigned BitSetOpcode = GL::FindILsb;
7039 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7041 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7043 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7046 return diagnoseUnsupported(
I,
7047 "spv_firstbitlow only supports 16,32,64 bits.");
7051bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
7052 SPIRVTypeInst ResType,
7053 MachineInstr &
I)
const {
7057 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
7060 .
addUse(
I.getOperand(2).getReg())
7063 unsigned Alignment =
I.getOperand(3).getImm();
7077 while (!Worklist.
empty()) {
7079 switch (
T->getOpcode()) {
7080 case SPIRV::OpTypeInt:
7081 case SPIRV::OpTypeFloat:
7082 case SPIRV::OpTypePointer:
7084 case SPIRV::OpTypeVector:
7085 case SPIRV::OpTypeVectorIdEXT:
7086 case SPIRV::OpTypeMatrix:
7087 case SPIRV::OpTypeArray: {
7088 Register OperandReg =
T->getOperand(1).getReg();
7092 case SPIRV::OpTypeStruct:
7093 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7094 Register OperandReg =
T->getOperand(Idx).getReg();
7109 Register TypeReg = Ty->getOperand(0).getReg();
7110 if (!Visited.
insert(TypeReg).second)
7113 switch (Ty->getOpcode()) {
7114 case SPIRV::OpTypePointer:
7115 if (Ty->getOperand(1).getImm() == SPIRV::StorageClass::StorageBuffer)
7119 case SPIRV::OpTypeArray:
7120 case SPIRV::OpTypeRuntimeArray:
7123 case SPIRV::OpTypeStruct:
7124 for (
unsigned I = 1;
I < Ty->getNumOperands(); ++
I)
7140bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7141 assert(
I.getNumExplicitOperands() == 2);
7143 Register MsgReg =
I.getOperand(1).getReg();
7145 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7148 return diagnoseUnsupported(
7150 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7151 "scalar, pointer, vector, matrix, or aggregate of such types)");
7154 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7161bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7170 uint32_t MsgVal = ~0
u;
7171 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7172 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7175 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7178 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7185bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7186 SPIRVTypeInst ResType,
7187 MachineInstr &
I)
const {
7194 bool UseUntypedPointers =
7195 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7197 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7200 MachineIRBuilder MIRBuilder(
I);
7203 .
addImm(SPIRV::Extension::SPV_KHR_variable_pointers);
7205 .
addImm(SPIRV::Capability::VariablePointersStorageBuffer);
7208 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7211 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7215 if (UseUntypedPointers) {
7219 return diagnoseUnsupported(
7220 I,
"could not deduce the data type of an untyped variable");
7226 unsigned Alignment =
I.getOperand(2).getImm();
7233bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7238 const MachineInstr *PrevI =
I.getPrevNode();
7240 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7244 .
addMBB(
I.getOperand(0).getMBB())
7249 .
addMBB(
I.getOperand(0).getMBB())
7254bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7265 const MachineInstr *NextI =
I.getNextNode();
7267 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7273 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7275 .
addUse(
I.getOperand(0).getReg())
7276 .
addMBB(
I.getOperand(1).getMBB())
7282bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7283 MachineInstr &
I)
const {
7285 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7287 const unsigned NumOps =
I.getNumOperands();
7288 for (
unsigned i = 1; i <
NumOps; i += 2) {
7289 MIB.
addUse(
I.getOperand(i + 0).getReg());
7290 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7296bool SPIRVInstructionSelector::selectGlobalValue(
7297 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7299 MachineIRBuilder MIRBuilder(
I);
7300 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7303 std::string GlobalIdent;
7305 unsigned &
ID = UnnamedGlobalIDs[GV];
7307 ID = UnnamedGlobalIDs.
size();
7308 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7334 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7341 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7346 MachineInstrBuilder MIB1 =
7347 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7350 MachineInstrBuilder MIB2 =
7352 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7356 GR.
add(ConstVal, MIB2);
7364 MachineInstrBuilder MIB3 =
7365 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7368 GR.
add(ConstVal, MIB3);
7374 assert(NewReg != ResVReg);
7375 return BuildCOPY(ResVReg, NewReg,
I);
7385 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7388 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7394 SPIRVTypeInst ResType =
7398 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7403 if (
GlobalVar->isExternallyInitialized() &&
7404 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7405 constexpr unsigned ReadWriteINTEL = 3u;
7408 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7414bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7415 SPIRVTypeInst ResType,
7416 MachineInstr &
I)
const {
7418 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7426 MachineIRBuilder MIRBuilder(
I);
7431 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7434 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7436 .
add(
I.getOperand(1))
7450 APFloat::rmNearestTiesToEven, &LosesInfo);
7455 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
7465bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7466 SPIRVTypeInst ResType,
7467 MachineInstr &
I)
const {
7470 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7476 Register ExpReg =
I.getOperand(2).getReg();
7478 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7479 SPIRV::OpConvertSToF))
7481 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7488bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7489 SPIRVTypeInst ResType,
7490 MachineInstr &
I)
const {
7506 MachineIRBuilder MIRBuilder(
I);
7507 SPIRVTypeInst FloatType =
7511 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7524 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7525 const bool IsUntyped =
7526 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7528 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7529 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7530 : SPIRV::OpVariable))
7533 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7541 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7544 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7547 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7551 Register IntegralPartReg =
I.getOperand(1).getReg();
7554 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7564 assert(
false &&
"GLSL::Modf is deprecated.");
7575bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7576 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7577 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7578 MachineIRBuilder MIRBuilder(
I);
7579 const SPIRVTypeInst Vec3Ty =
7582 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7594 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7598 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7604 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7611 assert(
I.getOperand(2).isReg());
7612 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7616 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7627bool SPIRVInstructionSelector::loadBuiltinInputID(
7628 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7629 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7630 MachineIRBuilder MIRBuilder(
I);
7632 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7647 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7651 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7660SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7661 MachineInstr &
I)
const {
7662 MachineIRBuilder MIRBuilder(
I);
7673bool SPIRVInstructionSelector::loadHandleBeforePosition(
7674 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7675 MachineInstr &Pos)
const {
7678 Intrinsic::spv_resource_handlefrombinding);
7686 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7687 MachineIRBuilder MIRBuilder(HandleDef);
7688 SPIRVTypeInst VarType = ResType;
7689 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7691 if (IsStructuredBuffer) {
7700 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7703 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7704 ArraySize, IndexReg, Name, MIRBuilder);
7708 uint32_t LoadOpcode =
7709 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7719bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7720 MachineInstr &
I)
const {
7722 return diagnoseUnsupported(
7723 I,
"this instruction is only supported in shaders.");
7728InstructionSelector *
7732 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