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;
173 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
182 Register castPtrToPtrToInt(
Register Ptr, SPIRV::StorageClass::StorageClass SC,
186 bool selectAtomicPtrValue(
202 unsigned OpType)
const;
270 unsigned Opcode)
const;
274 unsigned Opcode)
const;
278 unsigned Opcode)
const;
282 unsigned Opcode)
const;
284 template <
bool Signed>
287 template <
bool Signed>
294 template <
typename PickOpcodeFn>
297 PickOpcodeFn &&PickOpcode)
const;
314 template <
typename PickOpcodeFn>
317 PickOpcodeFn &&PickOpcode)
const;
335 bool IsSigned)
const;
337 bool IsSigned,
unsigned Opcode)
const;
339 bool IsSigned)
const;
345 bool IsSigned)
const;
386 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
387 bool useMISrc =
true,
389 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
390 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
391 bool useMISrc =
true,
393 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
394 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
395 bool setMIFlags =
true,
bool useMISrc =
true,
397 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
398 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
399 bool useMISrc =
true,
402 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
403 MachineInstr &
I)
const;
405 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
406 MachineInstr &
I)
const;
408 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
409 MachineInstr &
I)
const;
411 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
412 MachineInstr &
I,
unsigned Opcode)
const;
414 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
415 bool WithGroupSync)
const;
417 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
418 MachineInstr &
I)
const;
420 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
421 MachineInstr &
I)
const;
425 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
426 MachineInstr &
I)
const;
428 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
429 MachineInstr &
I)
const;
431 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
432 MachineInstr &
I)
const;
433 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
434 MachineInstr &
I)
const;
435 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
436 SPIRVTypeInst ResType,
437 MachineInstr &
I)
const;
438 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
439 MachineInstr &
I)
const;
442 std::optional<Register> LodReg = std::nullopt)
const;
443 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
444 MachineInstr &
I)
const;
445 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
446 MachineInstr &
I)
const;
447 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
450 MachineInstr &
I)
const;
451 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
452 MachineInstr &
I)
const;
453 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
454 MachineInstr &
I)
const;
455 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
458 SPIRVTypeInst ResType,
459 MachineInstr &
I)
const;
460 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
461 MachineInstr &
I)
const;
462 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
463 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
464 MachineInstr &
I)
const;
465 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
466 MachineInstr &
I)
const;
467 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
468 MachineInstr &
I)
const;
469 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
470 MachineInstr &
I)
const;
471 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
472 MachineInstr &
I)
const;
473 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
474 MachineInstr &
I)
const;
476 bool selectCopySign(
Register ResVReg, SPIRVTypeInst ResType,
477 MachineInstr &
I)
const;
479 bool selectLdexp(
Register ResVReg, SPIRVTypeInst ResType,
480 MachineInstr &
I)
const;
481 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
482 MachineInstr &
I)
const;
483 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
484 MachineInstr &
I)
const;
485 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
486 MachineInstr &
I,
const unsigned DPdOpCode)
const;
488 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
489 SPIRVTypeInst ResType =
nullptr)
const;
491 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
492 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
493 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
495 MachineInstr &
I)
const;
496 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
498 MachineInstr &
I)
const;
500 bool wrapIntoSpecConstantOp(MachineInstr &
I,
503 Register getUcharPtrTypeReg(MachineInstr &
I,
504 SPIRV::StorageClass::StorageClass SC)
const;
505 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
507 uint32_t Opcode)
const;
508 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
509 SPIRVTypeInst SrcPtrTy)
const;
510 Register buildPointerToResource(SPIRVTypeInst ResType,
511 SPIRV::StorageClass::StorageClass SC,
512 uint32_t Set, uint32_t
Binding,
513 uint32_t ArraySize,
Register IndexReg,
515 MachineIRBuilder MIRBuilder)
const;
516 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
517 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
518 Register &ReadReg, MachineInstr &InsertionPoint)
const;
519 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
522 const ImageOperands *ImOps =
nullptr)
const;
523 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
525 Register CoordinateReg,
const ImageOperands &ImOps,
528 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
529 Register ResVReg, SPIRVTypeInst ResType,
530 MachineInstr &
I)
const;
531 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
532 Register ResVReg, SPIRVTypeInst ResType,
533 MachineInstr &
I)
const;
534 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
535 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
536 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
537 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
540 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
541 Register SrcReg,
unsigned int Opcode,
542 std::function<
bool(
Register, SPIRVTypeInst,
543 MachineInstr &,
Register,
unsigned)>
547bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
549 if (
TET->getTargetExtName() ==
"spirv.Image") {
552 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
553 return TET->getTypeParameter(0)->isIntegerTy();
557#define GET_GLOBALISEL_IMPL
558#include "SPIRVGenGlobalISel.inc"
559#undef GET_GLOBALISEL_IMPL
565 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
568#include
"SPIRVGenGlobalISel.inc"
571#include
"SPIRVGenGlobalISel.inc"
583 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
588 if (HasVRegsReset == &MF)
603 for (
const auto &
MBB : MF) {
604 for (
const auto &
MI :
MBB) {
607 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
614 if (DstType != SrcType)
619 if (DstRC != SrcRC && SrcRC)
631 while (!Stack.empty()) {
636 switch (
MI->getOpcode()) {
637 case TargetOpcode::G_INTRINSIC:
638 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
639 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
642 if (IntrID != Intrinsic::spv_const_composite &&
643 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
647 case TargetOpcode::G_BUILD_VECTOR:
648 case TargetOpcode::G_SPLAT_VECTOR:
650 i < OpDef->getNumOperands(); i++) {
655 Stack.push_back(OpNestedDef);
658 case TargetOpcode::G_CONSTANT:
659 case TargetOpcode::G_FCONSTANT:
660 case TargetOpcode::G_IMPLICIT_DEF:
661 case SPIRV::OpConstantTrue:
662 case SPIRV::OpConstantFalse:
663 case SPIRV::OpConstantI:
664 case SPIRV::OpConstantF:
665 case SPIRV::OpConstantComposite:
666 case SPIRV::OpConstantCompositeContinuedINTEL:
667 case SPIRV::OpConstantSampler:
668 case SPIRV::OpConstantNull:
670 case SPIRV::OpPoisonKHR:
671 case SPIRV::OpConstantFunctionPointerINTEL:
698 case Intrinsic::spv_all:
699 case Intrinsic::spv_alloca:
700 case Intrinsic::spv_any:
701 case Intrinsic::spv_bitcast:
702 case Intrinsic::spv_const_composite:
703 case Intrinsic::spv_degrees:
704 case Intrinsic::spv_distance:
705 case Intrinsic::spv_extractelt:
706 case Intrinsic::spv_extractv:
707 case Intrinsic::spv_faceforward:
708 case Intrinsic::spv_fdot:
709 case Intrinsic::spv_firstbitlow:
710 case Intrinsic::spv_firstbitshigh:
711 case Intrinsic::spv_firstbituhigh:
712 case Intrinsic::spv_frac:
713 case Intrinsic::spv_gep:
714 case Intrinsic::spv_global_offset:
715 case Intrinsic::spv_global_size:
716 case Intrinsic::spv_group_id:
717 case Intrinsic::spv_insertelt:
718 case Intrinsic::spv_insertv:
719 case Intrinsic::spv_isinf:
720 case Intrinsic::spv_isnan:
721 case Intrinsic::spv_isfinite:
722 case Intrinsic::spv_isnormal:
723 case Intrinsic::spv_lerp:
724 case Intrinsic::spv_length:
725 case Intrinsic::spv_normalize:
726 case Intrinsic::spv_num_subgroups:
727 case Intrinsic::spv_num_workgroups:
728 case Intrinsic::spv_ptrcast:
729 case Intrinsic::spv_radians:
730 case Intrinsic::spv_reflect:
731 case Intrinsic::spv_refract:
732 case Intrinsic::spv_resource_getbasepointer:
733 case Intrinsic::spv_resource_getpointer:
734 case Intrinsic::spv_resource_handlefrombinding:
735 case Intrinsic::spv_resource_handlefromimplicitbinding:
736 case Intrinsic::spv_resource_nonuniformindex:
737 case Intrinsic::spv_resource_sample:
738 case Intrinsic::spv_rsqrt:
739 case Intrinsic::spv_saturate:
740 case Intrinsic::spv_sdot:
741 case Intrinsic::spv_sign:
742 case Intrinsic::spv_smoothstep:
743 case Intrinsic::spv_subgroup_id:
744 case Intrinsic::spv_subgroup_local_invocation_id:
745 case Intrinsic::spv_subgroup_max_size:
746 case Intrinsic::spv_subgroup_size:
747 case Intrinsic::spv_thread_id:
748 case Intrinsic::spv_thread_id_in_group:
749 case Intrinsic::spv_udot:
750 case Intrinsic::spv_undef:
751 case Intrinsic::spv_value_md:
752 case Intrinsic::spv_workgroup_size:
764 case SPIRV::OpTypeVoid:
765 case SPIRV::OpTypeBool:
766 case SPIRV::OpTypeInt:
767 case SPIRV::OpTypeFloat:
768 case SPIRV::OpTypeVector:
769 case SPIRV::OpTypeVectorIdEXT:
770 case SPIRV::OpTypeMatrix:
771 case SPIRV::OpTypeImage:
772 case SPIRV::OpTypeSampler:
773 case SPIRV::OpTypeSampledImage:
774 case SPIRV::OpTypeArray:
775 case SPIRV::OpTypeRuntimeArray:
776 case SPIRV::OpTypeStruct:
777 case SPIRV::OpTypeOpaque:
778 case SPIRV::OpTypePointer:
779 case SPIRV::OpTypeFunction:
780 case SPIRV::OpTypeEvent:
781 case SPIRV::OpTypeDeviceEvent:
782 case SPIRV::OpTypeReserveId:
783 case SPIRV::OpTypeQueue:
784 case SPIRV::OpTypePipe:
785 case SPIRV::OpTypeForwardPointer:
786 case SPIRV::OpTypePipeStorage:
787 case SPIRV::OpTypeNamedBarrier:
788 case SPIRV::OpTypeAccelerationStructureNV:
789 case SPIRV::OpTypeCooperativeMatrixNV:
790 case SPIRV::OpTypeCooperativeMatrixKHR:
800 if (
MI.getNumDefs() == 0)
803 for (
const auto &MO :
MI.all_defs()) {
805 if (
Reg.isPhysical()) {
809 for (
const auto &
UseMI :
MRI.use_nodbg_instructions(
Reg)) {
810 if (
UseMI.getOpcode() != SPIRV::OpName) {
817 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
818 MI.isLifetimeMarker()) {
821 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
832 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
833 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
836 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
841 if (
MI.mayStore() ||
MI.isCall() ||
842 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
843 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
844 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
855 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
862void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
864 for (
const auto &MO :
MI.all_defs()) {
868 SmallVector<MachineInstr *, 4> UselessOpNames;
871 "There is still a use of the dead function.");
874 for (MachineInstr *OpNameMI : UselessOpNames) {
876 OpNameMI->eraseFromParent();
881void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
884 removeOpNamesForDeadMI(
MI);
885 MI.eraseFromParent();
888bool SPIRVInstructionSelector::select(MachineInstr &
I) {
889 resetVRegsType(*
I.getParent()->getParent());
891 assert(
I.getParent() &&
"Instruction should be in a basic block!");
892 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
897 removeDeadInstruction(
I);
904 if (Opcode == SPIRV::ASSIGN_TYPE) {
905 Register DstReg =
I.getOperand(0).getReg();
906 Register SrcReg =
I.getOperand(1).getReg();
909 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
910 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
911 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
912 Register SelectDstReg =
Def->getOperand(0).getReg();
913 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
915 assert(SuccessToSelectSelect);
917 Def->eraseFromParent();
924 bool Res = selectImpl(
I, *CoverageInfo);
926 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
927 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
931 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
943 }
else if (
I.getNumDefs() == 1) {
955 removeDeadInstruction(
I);
960 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
961 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
967 bool HasDefs =
I.getNumDefs() > 0;
970 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
971 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
972 if (spvSelect(ResVReg, ResType,
I)) {
974 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
985 case TargetOpcode::G_CONSTANT:
986 case TargetOpcode::G_FCONSTANT:
993 MachineInstr &
I)
const {
996 if (DstRC != SrcRC && SrcRC)
998 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
1005bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
1006 SPIRVTypeInst ResType,
1007 MachineInstr &
I)
const {
1008 const unsigned Opcode =
I.getOpcode();
1010 return selectImpl(
I, *CoverageInfo);
1012 case TargetOpcode::G_CONSTANT:
1013 case TargetOpcode::G_FCONSTANT:
1014 return selectConst(ResVReg, ResType,
I);
1015 case TargetOpcode::G_GLOBAL_VALUE:
1016 return selectGlobalValue(ResVReg,
I);
1017 case TargetOpcode::G_IMPLICIT_DEF:
1018 return selectOpUndef(ResVReg, ResType,
I);
1019 case TargetOpcode::G_FREEZE:
1020 return selectFreeze(ResVReg, ResType,
I);
1022 case TargetOpcode::G_INTRINSIC:
1023 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1024 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1025 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1026 return selectIntrinsic(ResVReg, ResType,
I);
1027 case TargetOpcode::G_BITREVERSE:
1028 return selectBitreverse(ResVReg, ResType,
I);
1030 case TargetOpcode::G_BUILD_VECTOR:
1031 return selectBuildVector(ResVReg, ResType,
I);
1032 case TargetOpcode::G_SPLAT_VECTOR:
1033 return selectSplatVector(ResVReg, ResType,
I);
1034 case TargetOpcode::G_CONCAT_VECTORS:
1035 return selectConcatVectors(ResVReg, ResType,
I);
1037 case TargetOpcode::G_SHUFFLE_VECTOR: {
1038 MachineBasicBlock &BB = *
I.getParent();
1039 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1042 .
addUse(
I.getOperand(1).getReg())
1043 .
addUse(
I.getOperand(2).getReg());
1044 for (
auto V :
I.getOperand(3).getShuffleMask())
1049 case TargetOpcode::G_MEMMOVE:
1050 case TargetOpcode::G_MEMCPY:
1051 case TargetOpcode::G_MEMCPY_INLINE:
1052 case TargetOpcode::G_MEMSET:
1053 case TargetOpcode::G_MEMSET_INLINE:
1054 return selectMemOperation(ResVReg,
I);
1056 case TargetOpcode::G_ICMP:
1057 return selectICmp(ResVReg, ResType,
I);
1058 case TargetOpcode::G_FCMP:
1059 return selectFCmp(ResVReg, ResType,
I);
1061 case TargetOpcode::G_FRAME_INDEX:
1062 return selectFrameIndex(ResVReg, ResType,
I);
1064 case TargetOpcode::G_LOAD:
1065 return selectLoad(ResVReg, ResType,
I);
1066 case TargetOpcode::G_STORE:
1067 return selectStore(
I);
1069 case TargetOpcode::G_BR:
1070 return selectBranch(
I);
1071 case TargetOpcode::G_BRCOND:
1072 return selectBranchCond(
I);
1074 case TargetOpcode::G_PHI:
1075 return selectPhi(ResVReg,
I);
1077 case TargetOpcode::G_FPTOSI:
1078 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1079 case TargetOpcode::G_FPTOUI:
1080 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1082 case TargetOpcode::G_FPTOSI_SAT:
1083 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1084 case TargetOpcode::G_FPTOUI_SAT:
1085 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1087 case TargetOpcode::G_SITOFP:
1088 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1089 case TargetOpcode::G_UITOFP:
1090 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1092 case TargetOpcode::G_CTPOP:
1093 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1094 case TargetOpcode::G_SMIN:
1095 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1096 case TargetOpcode::G_UMIN:
1097 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1099 case TargetOpcode::G_SMAX:
1100 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1101 case TargetOpcode::G_UMAX:
1102 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1104 case TargetOpcode::G_SCMP:
1105 return selectSUCmp(ResVReg, ResType,
I,
true);
1106 case TargetOpcode::G_UCMP:
1107 return selectSUCmp(ResVReg, ResType,
I,
false);
1108 case TargetOpcode::G_LROUND:
1109 case TargetOpcode::G_LLROUND: {
1114 regForLround, *(
I.getParent()->getParent()));
1116 CL::round, GL::Round,
false);
1118 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1125 case TargetOpcode::G_STRICT_FMA:
1126 case TargetOpcode::G_FMA: {
1129 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1132 .
addUse(
I.getOperand(1).getReg())
1133 .
addUse(
I.getOperand(2).getReg())
1134 .
addUse(
I.getOperand(3).getReg())
1139 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1142 case TargetOpcode::G_FLDEXP:
1143 case TargetOpcode::G_STRICT_FLDEXP:
1144 return selectLdexp(ResVReg, ResType,
I);
1146 case TargetOpcode::G_FPOW:
1147 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1148 case TargetOpcode::G_FPOWI:
1149 return selectFpowi(ResVReg, ResType,
I);
1151 case TargetOpcode::G_FEXP:
1152 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1153 case TargetOpcode::G_FEXP2:
1154 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1155 case TargetOpcode::G_FEXP10:
1156 return selectExp10(ResVReg, ResType,
I);
1158 case TargetOpcode::G_FMODF:
1159 return selectModf(ResVReg, ResType,
I);
1160 case TargetOpcode::G_FSINCOS:
1161 return selectSincos(ResVReg, ResType,
I);
1163 case TargetOpcode::G_FLOG:
1164 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1165 case TargetOpcode::G_FLOG2:
1166 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1167 case TargetOpcode::G_FLOG10:
1168 return selectLog10(ResVReg, ResType,
I);
1170 case TargetOpcode::G_FABS:
1171 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1172 case TargetOpcode::G_ABS:
1173 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1175 case TargetOpcode::G_FMINNUM:
1176 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1177 case TargetOpcode::G_FMAXNUM:
1178 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1180 case TargetOpcode::G_FCOPYSIGN:
1181 return selectCopySign(ResVReg, ResType,
I);
1183 case TargetOpcode::G_FCEIL:
1184 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1185 case TargetOpcode::G_FFLOOR:
1186 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1188 case TargetOpcode::G_FCOS:
1189 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1190 case TargetOpcode::G_FSIN:
1191 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1192 case TargetOpcode::G_FTAN:
1193 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1194 case TargetOpcode::G_FACOS:
1195 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1196 case TargetOpcode::G_FASIN:
1197 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1198 case TargetOpcode::G_FATAN:
1199 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1200 case TargetOpcode::G_FATAN2:
1201 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1202 case TargetOpcode::G_FCOSH:
1203 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1204 case TargetOpcode::G_FSINH:
1205 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1206 case TargetOpcode::G_FTANH:
1207 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1209 case TargetOpcode::G_STRICT_FSQRT:
1210 case TargetOpcode::G_FSQRT:
1211 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1213 case TargetOpcode::G_CTTZ:
1214 case TargetOpcode::G_CTTZ_ZERO_POISON:
1215 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1216 case TargetOpcode::G_CTLZ:
1217 case TargetOpcode::G_CTLZ_ZERO_POISON:
1218 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1220 case TargetOpcode::G_INTRINSIC_ROUND:
1221 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1222 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1223 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1224 case TargetOpcode::G_INTRINSIC_TRUNC:
1225 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1226 case TargetOpcode::G_FRINT:
1227 case TargetOpcode::G_FNEARBYINT:
1228 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1230 case TargetOpcode::G_SMULH:
1231 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1232 case TargetOpcode::G_UMULH:
1233 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1235 case TargetOpcode::G_SADDSAT:
1236 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1237 case TargetOpcode::G_UADDSAT:
1238 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1239 case TargetOpcode::G_SSUBSAT:
1240 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1241 case TargetOpcode::G_USUBSAT:
1242 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1244 case TargetOpcode::G_FFREXP:
1245 return selectFrexp(ResVReg, ResType,
I);
1247 case TargetOpcode::G_UADDO:
1248 return selectOverflowArith(ResVReg, ResType,
I,
1250 : SPIRV::OpIAddCarryS);
1251 case TargetOpcode::G_USUBO:
1252 return selectOverflowArith(ResVReg, ResType,
I,
1254 : SPIRV::OpISubBorrowS);
1255 case TargetOpcode::G_UMULO:
1256 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1257 case TargetOpcode::G_SMULO:
1258 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1260 case TargetOpcode::G_SEXT:
1261 return selectExt(ResVReg, ResType,
I,
true);
1262 case TargetOpcode::G_ANYEXT:
1263 case TargetOpcode::G_ZEXT:
1264 return selectExt(ResVReg, ResType,
I,
false);
1265 case TargetOpcode::G_TRUNC:
1266 return selectTrunc(ResVReg, ResType,
I);
1267 case TargetOpcode::G_FPTRUNC:
1268 case TargetOpcode::G_FPEXT:
1269 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1271 case TargetOpcode::G_PTRTOINT:
1272 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1273 case TargetOpcode::G_INTTOPTR:
1274 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1275 case TargetOpcode::G_BITCAST:
1276 return selectBitcast(ResVReg, ResType,
I);
1277 case TargetOpcode::G_ADDRSPACE_CAST:
1278 return selectAddrSpaceCast(ResVReg, ResType,
I);
1279 case TargetOpcode::G_PTRMASK:
1280 return selectPtrMask(ResVReg, ResType,
I);
1281 case TargetOpcode::G_PTR_ADD: {
1283 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1287 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1288 (*II).getOpcode() == TargetOpcode::COPY ||
1289 (*II).getOpcode() == SPIRV::OpVariable ||
1290 (*II).getOpcode() == SPIRV::OpUntypedVariableKHR) &&
1293 bool IsGVInit =
false;
1297 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1298 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1299 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1300 (*UseIt).getOpcode() == SPIRV::OpVariable ||
1301 (*UseIt).getOpcode() == SPIRV::OpUntypedVariableKHR) {
1313 const bool UseUntypedPointers =
1314 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1315 if (UseUntypedPointers) {
1316 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1319 .
addImm(
static_cast<uint32_t
>(
1320 SPIRV::Opcode::UntypedInBoundsPtrAccessChainKHR))
1323 .
addUse(
I.getOperand(2).getReg())
1330 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1342 return diagnoseUnsupported(
1343 I,
"incompatible result and operand types in a bitcast");
1345 MachineInstrBuilder MIB =
1346 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1353 : SPIRV::OpInBoundsPtrAccessChain))
1357 .
addUse(
I.getOperand(2).getReg())
1360 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1364 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1366 .
addUse(
I.getOperand(2).getReg())
1375 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1378 .
addImm(
static_cast<uint32_t
>(
1379 SPIRV::Opcode::InBoundsPtrAccessChain))
1382 .
addUse(
I.getOperand(2).getReg());
1387 case TargetOpcode::G_ATOMICRMW_OR:
1388 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1389 case TargetOpcode::G_ATOMICRMW_ADD:
1390 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1391 case TargetOpcode::G_ATOMICRMW_AND:
1392 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1393 case TargetOpcode::G_ATOMICRMW_MAX:
1394 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1395 case TargetOpcode::G_ATOMICRMW_MIN:
1396 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1397 case TargetOpcode::G_ATOMICRMW_SUB:
1398 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1399 case TargetOpcode::G_ATOMICRMW_XOR:
1400 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1401 case TargetOpcode::G_ATOMICRMW_UMAX:
1402 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1403 case TargetOpcode::G_ATOMICRMW_UMIN:
1404 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1405 case TargetOpcode::G_ATOMICRMW_XCHG:
1406 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1408 case TargetOpcode::G_ATOMICRMW_FADD:
1409 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1410 case TargetOpcode::G_ATOMICRMW_FSUB:
1412 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1414 : SPIRV::OpFNegate);
1415 case TargetOpcode::G_ATOMICRMW_FMIN:
1416 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1417 case TargetOpcode::G_ATOMICRMW_FMAX:
1418 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1420 case TargetOpcode::G_FENCE:
1421 return selectFence(
I);
1423 case TargetOpcode::G_STACKSAVE:
1424 return selectStackSave(ResVReg, ResType,
I);
1425 case TargetOpcode::G_STACKRESTORE:
1426 return selectStackRestore(
I);
1428 case TargetOpcode::G_UNMERGE_VALUES:
1431 case TargetOpcode::G_TRAP:
1432 case TargetOpcode::G_UBSANTRAP:
1433 return selectTrap(
I);
1438 case TargetOpcode::DBG_LABEL:
1440 case TargetOpcode::G_DEBUGTRAP:
1441 return selectDebugTrap(ResVReg, ResType,
I);
1442 case TargetOpcode::G_PREFETCH:
1443 return selectPrefetch(
I);
1450bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1451 SPIRVTypeInst ResType,
1452 MachineInstr &
I)
const {
1453 unsigned Opcode = SPIRV::OpNop;
1460bool SPIRVInstructionSelector::selectPrefetch(MachineInstr &
I)
const {
1469 MachineIRBuilder MIRBuilder(
I);
1471 const SPIRVTypeInst PointerSizeType =
1479 Register AddrVal =
I.getOperand(0).getReg();
1482 return selectExtInst(ExtReg, GR.
getOpTypeVoid(MIRBuilder),
I, CL::prefetch,
1484 {AddrVal, ConstIntOne});
1489bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1490 SPIRVTypeInst ResType,
1492 GL::GLSLExtInst GLInst,
1493 bool setMIFlags,
bool useMISrc,
1496 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1497 return diagnoseUnsupported(
1499 "this instruction is only supported with the GLSL extended instruction "
1501 return selectExtInst(ResVReg, ResType,
I,
1502 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1503 setMIFlags, useMISrc, SrcRegs);
1506bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1507 SPIRVTypeInst ResType,
1509 CL::OpenCLExtInst CLInst,
1510 bool setMIFlags,
bool useMISrc,
1512 return selectExtInst(ResVReg, ResType,
I,
1513 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1514 setMIFlags, useMISrc, SrcRegs);
1517bool SPIRVInstructionSelector::selectExtInst(
1518 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1519 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1521 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1522 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1523 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1527bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1528 SPIRVTypeInst ResType,
1531 bool setMIFlags,
bool useMISrc,
1534 for (
const auto &[InstructionSet, Opcode] : Insts) {
1538 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1541 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1546 const unsigned NumOps =
I.getNumOperands();
1549 I.getOperand(Index).getType() ==
1550 MachineOperand::MachineOperandType::MO_IntrinsicID)
1553 MIB.
add(
I.getOperand(Index));
1565bool SPIRVInstructionSelector::selectCopySign(
Register ResVReg,
1566 SPIRVTypeInst ResType,
1567 MachineInstr &
I)
const {
1569 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1574 Register MagnitudeReg =
I.getOperand(1).getReg();
1575 Register SignReg =
I.getOperand(2).getReg();
1583 unsigned AndOpcode, OrOpcode;
1584 if (ComponentCount > 1) {
1588 AndOpcode = SPIRV::OpBitwiseAndV;
1589 OrOpcode = SPIRV::OpBitwiseOrV;
1593 AndOpcode = SPIRV::OpBitwiseAndS;
1594 OrOpcode = SPIRV::OpBitwiseOrS;
1600 return selectOpWithSrcs(ResReg, IntType,
I, SrcRegs, Opcode);
1603 Register MagnitudeInt, SignInt, MagnitudeBits, SignBits, CombinedInt;
1604 if (!EmitBitOp(MagnitudeInt, {MagnitudeReg}, SPIRV::OpBitcast) ||
1605 !EmitBitOp(SignInt, {SignReg}, SPIRV::OpBitcast) ||
1606 !EmitBitOp(MagnitudeBits, {MagnitudeInt, NotSignMask}, AndOpcode) ||
1607 !EmitBitOp(SignBits, {SignInt, SignMask}, AndOpcode) ||
1608 !EmitBitOp(CombinedInt, {MagnitudeBits, SignBits}, OrOpcode))
1611 return selectOpWithSrcs(ResVReg, ResType,
I, {CombinedInt}, SPIRV::OpBitcast);
1614bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1615 SPIRVTypeInst ResType,
1616 MachineInstr &
I)
const {
1617 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1618 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1619 for (
const auto &Ex : ExtInsts) {
1620 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1621 uint32_t Opcode = Ex.second;
1625 MachineIRBuilder MIRBuilder(
I);
1632 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1639 const bool IsUntyped =
1640 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1642 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1643 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1644 : SPIRV::OpVariable))
1647 .
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 ExpReg = buildVectorSplat(ExpReg, NumElts,
I);
1689 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1690 true,
false, {XReg, ExpReg});
1693bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1694 SPIRVTypeInst ResType,
1695 MachineInstr &
I)
const {
1696 Register CosResVReg =
I.getOperand(1).getReg();
1697 unsigned SrcIdx =
I.getNumExplicitDefs();
1702 MachineIRBuilder MIRBuilder(
I);
1704 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1711 const bool IsUntyped =
1712 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1714 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1715 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1716 : SPIRV::OpVariable))
1719 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1723 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1726 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1728 .
add(
I.getOperand(SrcIdx))
1732 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1740 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1743 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1745 .
add(
I.getOperand(SrcIdx))
1747 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1750 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1752 .
add(
I.getOperand(SrcIdx))
1759bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1760 SPIRVTypeInst ResType,
1763 unsigned Opcode)
const {
1764 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1774bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1775 SPIRVTypeInst ResType,
1778 unsigned Opcode)
const {
1779 MachineIRBuilder MIRBuilder(
I);
1781 Register OpReg =
I.getOperand(1).getReg();
1787 SPIRVTypeInst ExtType =
1794 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1798 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1801 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1804bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1805 SPIRVTypeInst ResType,
1808 unsigned Opcode)
const {
1809 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1812bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1813 SPIRVTypeInst ResType,
1816 unsigned Opcode)
const {
1817 MachineIRBuilder MIRBuilder(
I);
1826 SPIRVTypeInst WorkingType =
1833 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {SrcReg}, SPIRV::OpUConvert))
1837 if (!selectOpWithSrcs(LowCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1845 IsScalar ? SPIRV::OpShiftRightLogicalS : SPIRV::OpShiftRightLogicalV;
1847 if (!selectOpWithSrcs(Shift, SrcType,
I, {SrcReg, ShiftAmount}, ShiftOp))
1851 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {Shift}, SPIRV::OpUConvert))
1855 if (!selectOpWithSrcs(HighCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1860 if (!selectOpWithSrcs(Sum, WorkingType,
I, {HighCount, LowCount},
1861 IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV))
1865 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1866 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1869bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1870 SPIRVTypeInst ResType,
1872 unsigned Opcode)
const {
1877 if (!STI.getTargetTriple().isVulkanOS())
1878 return selectUnOp(ResVReg, ResType,
I, Opcode);
1880 Register OpReg =
I.getOperand(1).getReg();
1883 : SPIRV::OpUConvert;
1887 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1889 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1891 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1893 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1897bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1898 SPIRVTypeInst ResType,
1900 unsigned Opcode)
const {
1902 Register SrcReg =
I.getOperand(1).getReg();
1907 unsigned DefOpCode = DefIt->getOpcode();
1908 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1911 if (
auto *VRD =
getVRegDef(*
MRI, DefIt->getOperand(1).getReg()))
1912 DefOpCode = VRD->getOpcode();
1914 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1915 DefOpCode == TargetOpcode::G_CONSTANT ||
1916 DefOpCode == SPIRV::OpVariable ||
1917 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1918 DefOpCode == SPIRV::OpConstantI) {
1924 uint32_t SpecOpcode = 0;
1926 case SPIRV::OpConvertPtrToU:
1927 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1929 case SPIRV::OpConvertUToPtr:
1930 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1935 TII.get(SPIRV::OpSpecConstantOp))
1945 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1949bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1950 SPIRVTypeInst ResType,
1951 MachineInstr &
I)
const {
1952 Register OpReg =
I.getOperand(1).getReg();
1953 SPIRVTypeInst OpType =
1956 return diagnoseUnsupported(
1957 I,
"incompatible result and operand types in a bitcast");
1958 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1964 std::optional<Align> AlignOverride = std::nullopt) {
1969 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1970 if (
MemOp->isNonTemporal())
1971 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1973 if (!ST->isShader() &&
MemOp->getAlign().value())
1974 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1978 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1979 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1983 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1985 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1989 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1993 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1995 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1996 MIB.
addImm(AlignOverride.value_or(
MemOp->getAlign()).value());
2007 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2009 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2011 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2015bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2016 SPIRVTypeInst ResType,
2017 MachineInstr &
I)
const {
2019 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2024 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2025 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2027 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2029 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2033 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2037 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2038 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2039 I.getDebugLoc(),
I);
2043 MachineIRBuilder MIRBuilder(
I);
2045 if (
I.getNumMemOperands()) {
2046 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2047 if (MemOp->isAtomic())
2048 return selectAtomicLoad(ResVReg, ResType,
I);
2051 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2055 if (!
I.getNumMemOperands()) {
2056 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2058 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2067Register SPIRVInstructionSelector::createPtrSizedIntReg(
2068 MachineIRBuilder &MIRBuilder)
const {
2069 SPIRVTypeInst IntType =
2079SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2080 MachineIRBuilder &MIRBuilder)
const {
2081 SPIRVTypeInst IntType =
2083 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2084 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2092Register SPIRVInstructionSelector::castPtrToPtrToInt(
2093 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2094 MachineIRBuilder &MIRBuilder)
const {
2095 SPIRVTypeInst IntType =
2097 SPIRVTypeInst PtrType =
2111bool SPIRVInstructionSelector::selectAtomicPtrValue(
2112 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2113 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2121 Register IntResult = EmitAtomic(IntType);
2123 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2131bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2132 SPIRVTypeInst ResType,
2133 MachineInstr &
I)
const {
2134 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2137 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2140 return diagnoseUnsupported(
2141 I,
"Lowering to SPIR-V of atomic load is only "
2142 "allowed for integer, floating point or pointer types");
2144 assert(
I.getNumMemOperands());
2145 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2146 assert(MemOp.isAtomic());
2148 uint32_t
Scope =
static_cast<uint32_t
>(
2150 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2156 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2157 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2160 Register MemSemReg = buildI32ConstantInEntryBlock(Sem,
I);
2162 MachineIRBuilder MIRBuilder(
I);
2166 return diagnoseUnsupported(
2167 I,
"Lowering to SPIR-V of atomic load is only "
2168 "allowed for pointer types for physical addressing model");
2173 SPIRV::StorageClass::StorageClass SC =
2175 return selectAtomicPtrValue(
2176 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2177 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2178 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2189 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2200bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2202 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2203 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2208 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2209 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2211 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2216 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2220 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2221 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2222 SPIRVTypeInst SampledType =
2224 SPIRVTypeInst StoreValCompType =
2226 if (StoreValCompType && StoreValCompType != SampledType) {
2229 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2232 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2237 StoreVal = PackedReg;
2240 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2241 TII.get(SPIRV::OpImageWrite))
2247 if (sampledTypeIsSignedInteger(LLVMHandleType))
2250 BMI.constrainAllUses(
TII,
TRI, RBI);
2257 if (PointeeTy && PointeeTy->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
2258 StoreTy->
getOpcode() != SPIRV::OpTypeVectorIdEXT &&
2272 if (
I.getNumMemOperands()) {
2273 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2274 if (MemOp->isAtomic())
2275 return selectAtomicStore(
I);
2282 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2283 PtrSC == SPIRV::StorageClass::Input ||
2284 PtrSC == SPIRV::StorageClass::PushConstant)
2285 return diagnoseUnsupported(
2286 I,
"store into a read-only SPIR-V storage class is not allowed");
2288 MachineIRBuilder MIRBuilder(
I);
2290 if (!
I.getNumMemOperands()) {
2291 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2293 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2302bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2303 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2306 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2307 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2312 if (!PointeeType && PtrType &&
2313 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2316 return diagnoseUnsupported(
I,
2317 "Lowering to SPIR-V of atomic store is only "
2318 "allowed for integer or floating point types");
2320 assert(
I.getNumMemOperands());
2321 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2322 assert(MemOp.isAtomic());
2324 uint32_t
Scope =
static_cast<uint32_t
>(
2326 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2332 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2333 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2336 Register MemSemReg = buildI32ConstantInEntryBlock(Sem,
I);
2337 MachineIRBuilder MIRBuilder(
I);
2341 return diagnoseUnsupported(
2342 I,
"Lowering to SPIR-V of atomic store is only "
2343 "allowed for pointer types for physical addressing model");
2348 SPIRV::StorageClass::StorageClass SC =
2350 return selectAtomicPtrValue(
2351 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2353 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2366 return diagnoseUnsupported(
I,
2367 "Lowering to SPIR-V of atomic store is only "
2368 "allowed for integer or floating point types");
2370 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2380bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2381 SPIRVTypeInst ResType,
2382 MachineInstr &
I)
const {
2383 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2391 const Register PtrsReg =
I.getOperand(2).getReg();
2392 const uint32_t
Alignment =
I.getOperand(3).getImm();
2393 const Register MaskReg =
I.getOperand(4).getReg();
2394 const Register PassthruReg =
I.getOperand(5).getReg();
2395 const Register AlignmentReg = buildI32ConstantInEntryBlock(Alignment,
I);
2399 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2410bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2411 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2418 const Register ValuesReg =
I.getOperand(1).getReg();
2419 const Register PtrsReg =
I.getOperand(2).getReg();
2420 const uint32_t
Alignment =
I.getOperand(3).getImm();
2421 const Register MaskReg =
I.getOperand(4).getReg();
2422 const Register AlignmentReg = buildI32ConstantInEntryBlock(Alignment,
I);
2426 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2435bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2436 const Twine &
Msg)
const {
2437 const Function &
F =
I.getMF()->getFunction();
2438 F.getContext().diagnose(
2439 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2443bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2444 SPIRVTypeInst ResType,
2445 MachineInstr &
I)
const {
2446 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2447 return diagnoseUnsupported(
2448 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2449 "SPIR-V extension: SPV_INTEL_variable_length_array");
2451 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2458bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2459 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2460 return diagnoseUnsupported(
2462 "llvm.stackrestore intrinsic: this instruction requires the following "
2463 "SPIR-V extension: SPV_INTEL_variable_length_array");
2464 if (!
I.getOperand(0).isReg())
2467 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2468 .
addUse(
I.getOperand(0).getReg())
2474SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2475 MachineIRBuilder MIRBuilder(
I);
2476 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2483 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2487 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2488 Type *ArrTy = ArrayType::get(ValTy, Num);
2490 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2493 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2504 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2505 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2506 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2507 : SPIRV::OpVariable))
2510 .
addImm(SPIRV::StorageClass::UniformConstant);
2523bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2526 Register DstReg =
I.getOperand(0).getReg();
2532 return diagnoseUnsupported(
2533 I,
"OpCopyMemory requires operands to have the same type");
2538 return diagnoseUnsupported(
2539 I,
"Unable to determine pointee type size for OpCopyMemory");
2540 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2541 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2542 return diagnoseUnsupported(
2543 I,
"OpCopyMemory requires the size to match the pointee type size");
2546 const unsigned Opcode =
2547 IsLogical ? SPIRV::OpCopyMemory : SPIRV::OpCopyMemorySized;
2548 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2552 MIB.
addUse(
I.getOperand(2).getReg());
2554 if (
I.getNumMemOperands()) {
2555 MachineIRBuilder MIRBuilder(
I);
2556 const MachineMemOperand *DstMemOp = *
I.memoperands_begin();
2558 Align SrcAlign = DstAlign;
2561 if (
I.getNumMemOperands() > 1 && !STI.
isShader())
2562 SrcAlign = (*std::next(
I.memoperands_begin()))->
getAlign();
2572 std::min(DstAlign, SrcAlign));
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))
2601 if (!selectCopyMemory(
I, SrcReg))
2603 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2604 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2609bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2610 SPIRVTypeInst ResType,
2613 unsigned NegateOpcode)
const {
2615 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2616 uint32_t
Scope =
static_cast<uint32_t
>(
2618 MemOp->getSyncScopeID()));
2619 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2621 Register Ptr =
I.getOperand(1).getReg();
2622 uint32_t ScSem =
static_cast<uint32_t
>(
2626 Register MemSemReg = buildI32ConstantInEntryBlock(
2630 Register ValueReg =
I.getOperand(2).getReg();
2631 if (NegateOpcode != 0) {
2634 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2640 if (NewOpcode != SPIRV::OpAtomicExchange)
2641 return diagnoseUnsupported(
2642 I,
"Lowering to SPIR-V of this atomic operation is not "
2643 "allowed for pointer types");
2645 return diagnoseUnsupported(
2646 I,
"Lowering to SPIR-V of atomic exchange is only "
2647 "allowed for pointer types for physical addressing model");
2654 MachineIRBuilder MIRBuilder(
I);
2656 return selectAtomicPtrValue(
2657 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2659 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2660 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2661 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2669 return ExchangeResReg;
2673 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2684bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2685 unsigned ArgI =
I.getNumOperands() - 1;
2687 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2688 SPIRVTypeInst SrcType =
2692 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2696 unsigned CurrentIndex = 0;
2697 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2698 Register ResVReg =
I.getOperand(i).getReg();
2707 ResType = ScalarType;
2716 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2722 for (
unsigned j = 0;
j < NumElements; ++
j) {
2723 MIB.
addImm(CurrentIndex + j);
2725 CurrentIndex += NumElements;
2729 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2741bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2744 ? SPIRV::MemorySemantics::UniformMemory |
2745 SPIRV::MemorySemantics::WorkgroupMemory |
2746 SPIRV::MemorySemantics::ImageMemory
2747 : SPIRV::MemorySemantics::WorkgroupMemory |
2748 SPIRV::MemorySemantics::CrossWorkgroupMemory |
2749 SPIRV::MemorySemantics::ImageMemory;
2751 STI.getTargetTriple(),
static_cast<uint32_t
>(
getMemSemantics(AO)), ScSem);
2752 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
2756 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2758 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2765bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2766 SPIRVTypeInst ResType,
2768 unsigned Opcode)
const {
2769 Type *ResTy =
nullptr;
2772 return diagnoseUnsupported(
2774 "Not enough info to select the arithmetic with overflow instruction");
2776 return diagnoseUnsupported(
I,
2777 "Expect struct type result for the arithmetic "
2778 "with overflow instruction");
2784 MachineIRBuilder MIRBuilder(
I);
2786 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2787 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2794 Register ZeroReg = buildZerosVal(ResType,
I);
2799 if (ResName.
size() > 0)
2807 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2808 MIB.
addUse(
I.getOperand(i).getReg());
2814 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2816 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2817 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2824 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2825 .
addDef(
I.getOperand(1).getReg())
2833bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2834 SPIRVTypeInst ResType,
2835 MachineInstr &
I)
const {
2837 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2838 Register Ptr =
I.getOperand(2).getReg();
2839 Register ScopeReg =
I.getOperand(5).getReg();
2840 Register MemSemEqReg =
I.getOperand(6).getReg();
2841 Register MemSemNeqReg =
I.getOperand(7).getReg();
2843 Register Val =
I.getOperand(4).getReg();
2847 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2866 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2873 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2885 case SPIRV::StorageClass::DeviceOnlyINTEL:
2886 case SPIRV::StorageClass::HostOnlyINTEL:
2895 bool IsGRef =
false;
2896 bool IsAllowedRefs =
2897 llvm::all_of(
MRI->use_instructions(ResVReg), [&IsGRef](
auto const &It) {
2898 unsigned Opcode = It.getOpcode();
2899 if (Opcode == SPIRV::OpConstantComposite ||
2900 Opcode == SPIRV::OpSpecConstantComposite ||
2901 Opcode == SPIRV::OpVariable ||
2902 Opcode == SPIRV::OpUntypedVariableKHR ||
2903 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2904 return IsGRef = true;
2905 return Opcode == SPIRV::OpName;
2907 return IsAllowedRefs && IsGRef;
2910Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2911 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2913 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2917SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2919 uint32_t Opcode)
const {
2920 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2921 TII.get(SPIRV::OpSpecConstantOp))
2929SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2930 SPIRVTypeInst SrcPtrTy)
const {
2931 SPIRVTypeInst GenericPtrTy =
2935 SPIRV::StorageClass::Generic),
2939 MachineInstrBuilder MIB = buildSpecConstantOp(
2941 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2951bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2952 SPIRVTypeInst ResType,
2953 MachineInstr &
I)
const {
2957 Register SrcPtr =
I.getOperand(1).getReg();
2962 return BuildCOPY(ResVReg, SrcPtr,
I);
2972 unsigned SpecOpcode = [&]() ->
unsigned {
2973 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2974 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2976 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2978 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
2986 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2988 .constrainAllUses(
TII,
TRI, RBI);
2990 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2992 buildSpecConstantOp(
2994 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2995 .constrainAllUses(
TII,
TRI, RBI);
3002 return BuildCOPY(ResVReg, SrcPtr,
I);
3004 if ((SrcSC == SPIRV::StorageClass::Function &&
3005 DstSC == SPIRV::StorageClass::Private) ||
3006 (DstSC == SPIRV::StorageClass::Function &&
3007 SrcSC == SPIRV::StorageClass::Private))
3008 return BuildCOPY(ResVReg, SrcPtr,
I);
3012 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3015 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3018 SPIRVTypeInst GenericPtrTy =
3037 return selectUnOp(ResVReg, ResType,
I,
3038 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
3040 return selectUnOp(ResVReg, ResType,
I,
3041 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
3043 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3045 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3055bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3056 SPIRVTypeInst ResType,
3057 MachineInstr &
I)
const {
3059 return diagnoseUnsupported(
3060 I,
"G_PTRMASK is not supported with logical SPIR-V");
3065 Register PtrReg =
I.getOperand(1).getReg();
3066 Register MaskReg =
I.getOperand(2).getReg();
3085 ? SPIRV::OpBitwiseAndV
3086 : SPIRV::OpBitwiseAndS;
3109 return SPIRV::OpFOrdEqual;
3111 return SPIRV::OpFOrdGreaterThanEqual;
3113 return SPIRV::OpFOrdGreaterThan;
3115 return SPIRV::OpFOrdLessThanEqual;
3117 return SPIRV::OpFOrdLessThan;
3119 return SPIRV::OpFOrdNotEqual;
3121 return SPIRV::OpOrdered;
3123 return SPIRV::OpFUnordEqual;
3125 return SPIRV::OpFUnordGreaterThanEqual;
3127 return SPIRV::OpFUnordGreaterThan;
3129 return SPIRV::OpFUnordLessThanEqual;
3131 return SPIRV::OpFUnordLessThan;
3133 return SPIRV::OpFUnordNotEqual;
3135 return SPIRV::OpUnordered;
3145 return SPIRV::OpIEqual;
3147 return SPIRV::OpINotEqual;
3149 return SPIRV::OpSGreaterThanEqual;
3151 return SPIRV::OpSGreaterThan;
3153 return SPIRV::OpSLessThanEqual;
3155 return SPIRV::OpSLessThan;
3157 return SPIRV::OpUGreaterThanEqual;
3159 return SPIRV::OpUGreaterThan;
3161 return SPIRV::OpULessThanEqual;
3163 return SPIRV::OpULessThan;
3172 return SPIRV::OpPtrEqual;
3174 return SPIRV::OpPtrNotEqual;
3185 return SPIRV::OpLogicalEqual;
3187 return SPIRV::OpLogicalNotEqual;
3225bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3226 SPIRVTypeInst ResType,
3228 unsigned OpAnyOrAll)
const {
3229 assert(
I.getNumOperands() == 3);
3230 assert(
I.getOperand(2).isReg());
3232 Register InputRegister =
I.getOperand(2).getReg();
3235 assert(InputType &&
"VReg has no type assigned");
3239 assert(ResVReg ==
I.getOperand(0).getReg());
3240 return BuildCOPY(ResVReg, InputRegister,
I);
3244 unsigned SpirvNotEqualId =
3245 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3247 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3252 IsBoolTy ? InputRegister
3260 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3262 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3279bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3280 SPIRVTypeInst ResType,
3281 MachineInstr &
I)
const {
3282 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3285bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3286 SPIRVTypeInst ResType,
3287 MachineInstr &
I)
const {
3288 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3292bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3293 SPIRVTypeInst ResType,
3294 MachineInstr &
I)
const {
3295 assert(
I.getNumOperands() == 4);
3296 assert(
I.getOperand(2).isReg());
3297 assert(
I.getOperand(3).isReg());
3299 [[maybe_unused]] SPIRVTypeInst VecType =
3304 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3305 "dot product requires either a vector of at least 2 components or"
3306 " the SPV_EXT_long vector extension.");
3308 [[maybe_unused]] SPIRVTypeInst EltType =
3317 .
addUse(
I.getOperand(2).getReg())
3318 .
addUse(
I.getOperand(3).getReg())
3323bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3324 SPIRVTypeInst ResType,
3327 assert(
I.getNumOperands() == 4);
3328 assert(
I.getOperand(2).isReg());
3329 assert(
I.getOperand(3).isReg());
3332 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3336 .
addUse(
I.getOperand(2).getReg())
3337 .
addUse(
I.getOperand(3).getReg())
3344bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3345 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3346 assert(
I.getNumOperands() == 4);
3347 assert(
I.getOperand(2).isReg());
3348 assert(
I.getOperand(3).isReg());
3352 Register Vec0 =
I.getOperand(2).getReg();
3353 Register Vec1 =
I.getOperand(3).getReg();
3357 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3366 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3367 "dot product requires either a vector of at least 2 components "
3368 "or the SPV_EXT_long_vector extension.");
3371 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3381 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3392 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3404bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3405 SPIRVTypeInst ResType,
3406 MachineInstr &
I)
const {
3408 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3411 .
addUse(
I.getOperand(2).getReg())
3416bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3417 SPIRVTypeInst ResType,
3418 MachineInstr &
I)
const {
3420 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3423 .
addUse(
I.getOperand(2).getReg())
3428bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3429 SPIRVTypeInst ResType,
3430 MachineInstr &
I)
const {
3433 Register Src =
I.getOperand(2).getReg();
3476bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3477 SPIRVTypeInst ResType,
3478 MachineInstr &
I)
const {
3480 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3483 .
addUse(
I.getOperand(2).getReg())
3488template <
bool Signed>
3489bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3490 SPIRVTypeInst ResType,
3491 MachineInstr &
I)
const {
3492 assert(
I.getNumOperands() == 5);
3493 assert(
I.getOperand(2).isReg());
3494 assert(
I.getOperand(3).isReg());
3495 assert(
I.getOperand(4).isReg());
3498 Register Acc =
I.getOperand(2).getReg();
3502 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3504 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3509 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3512 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3524template <
bool Signed>
3525bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3526 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3527 assert(
I.getNumOperands() == 5);
3528 assert(
I.getOperand(2).isReg());
3529 assert(
I.getOperand(3).isReg());
3530 assert(
I.getOperand(4).isReg());
3533 Register Acc =
I.getOperand(2).getReg();
3539 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3543 for (
unsigned i = 0; i < 4; i++) {
3566 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3586 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3601bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3602 SPIRVTypeInst ResType,
3603 MachineInstr &
I)
const {
3604 assert(
I.getNumOperands() == 3);
3605 assert(
I.getOperand(2).isReg());
3607 Register VZero = buildZerosValF(ResType,
I);
3608 Register VOne = buildOnesValF(ResType,
I);
3610 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3613 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3615 .
addUse(
I.getOperand(2).getReg())
3622bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3623 SPIRVTypeInst ResType,
3624 MachineInstr &
I)
const {
3625 assert(
I.getNumOperands() == 3);
3626 assert(
I.getOperand(2).isReg());
3628 Register InputRegister =
I.getOperand(2).getReg();
3630 auto &
DL =
I.getDebugLoc();
3633 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3640 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3642 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3650 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3655 if (NeedsConversion) {
3656 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3667bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3668 SPIRVTypeInst ResType,
3670 unsigned Opcode)
const {
3674 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3680 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3681 BMI.addUse(
I.getOperand(J).getReg());
3688bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3691 bool WithGroupSync)
const {
3693 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3695 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3697 assert(((Scope != SPIRV::Scope::Workgroup) ||
3698 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3699 "Workgroup Scope must set WorkGroupMemory semantic "
3700 "in Barrier instruction");
3702 assert(((Scope != SPIRV::Scope::Device) ||
3703 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3704 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3705 "Device Scope must set UniformMemory and ImageMemory semantic "
3706 "in Barrier instruction");
3712 if (WithGroupSync) {
3713 Register ExecReg = buildI32ConstantInEntryBlock(SPIRV::Scope::Workgroup,
I);
3717 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
3718 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
3720 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3724bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3725 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3730 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3731 SPIRV::OpGroupNonUniformBallot))
3736 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3741 .
addImm(SPIRV::GroupOperation::Reduce)
3748bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3749 SPIRVTypeInst ResType,
3750 MachineInstr &
I)
const {
3755 Register InputReg =
I.getOperand(2).getReg();
3760 bool IsVector = NumElems > 1 ||
3761 (InputType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
3775 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3776 SPIRV::OpGroupNonUniformAllEqual);
3781 ElementResults.
reserve(NumElems);
3783 for (
unsigned Idx = 0;
Idx < NumElems; ++
Idx) {
3796 ElemInput = Extracted;
3802 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3813 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3824bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3825 SPIRVTypeInst ResType,
3826 MachineInstr &
I)
const {
3828 assert(
I.getNumOperands() == 3);
3830 auto Op =
I.getOperand(2);
3840 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3842 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3843 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3864 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3868 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3875bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3876 SPIRVTypeInst ResType,
3878 bool IsUnsigned)
const {
3879 return selectWaveReduce(
3880 ResVReg, ResType,
I, IsUnsigned,
3881 [&](
Register InputRegister,
bool IsUnsigned) {
3882 const bool IsFloatTy =
3884 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3885 : SPIRV::OpGroupNonUniformSMax;
3886 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3890bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3891 SPIRVTypeInst ResType,
3893 bool IsUnsigned)
const {
3894 return selectWaveReduce(
3895 ResVReg, ResType,
I, IsUnsigned,
3896 [&](
Register InputRegister,
bool IsUnsigned) {
3897 const bool IsFloatTy =
3899 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3900 : SPIRV::OpGroupNonUniformSMin;
3901 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3905bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3906 SPIRVTypeInst ResType,
3907 MachineInstr &
I)
const {
3908 return selectWaveReduce(ResVReg, ResType,
I,
false,
3909 [&](
Register InputRegister,
bool IsUnsigned) {
3911 InputRegister, SPIRV::OpTypeFloat);
3912 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3913 : SPIRV::OpGroupNonUniformIAdd;
3917bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3918 SPIRVTypeInst ResType,
3919 MachineInstr &
I)
const {
3920 return selectWaveReduce(ResVReg, ResType,
I,
false,
3921 [&](
Register InputRegister,
bool IsUnsigned) {
3923 InputRegister, SPIRV::OpTypeFloat);
3924 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3925 : SPIRV::OpGroupNonUniformIMul;
3929template <
typename PickOpcodeFn>
3930bool SPIRVInstructionSelector::selectWaveReduce(
3931 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3932 PickOpcodeFn &&PickOpcode)
const {
3933 assert(
I.getNumOperands() == 3);
3934 assert(
I.getOperand(2).isReg());
3936 Register InputRegister =
I.getOperand(2).getReg();
3940 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3943 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3949 .
addImm(SPIRV::GroupOperation::Reduce)
3950 .
addUse(
I.getOperand(2).getReg())
3955bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3956 SPIRVTypeInst ResType,
3958 unsigned Opcode)
const {
3959 return selectWaveReduce(
3960 ResVReg, ResType,
I,
false,
3961 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3964bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3965 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3966 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3967 [&](
Register InputRegister,
bool IsUnsigned) {
3969 InputRegister, SPIRV::OpTypeFloat);
3971 ? SPIRV::OpGroupNonUniformFAdd
3972 : SPIRV::OpGroupNonUniformIAdd;
3976bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3977 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3978 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3979 [&](
Register InputRegister,
bool IsUnsigned) {
3981 InputRegister, SPIRV::OpTypeFloat);
3983 ? SPIRV::OpGroupNonUniformFMul
3984 : SPIRV::OpGroupNonUniformIMul;
3988template <
typename PickOpcodeFn>
3989bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3990 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3991 PickOpcodeFn &&PickOpcode)
const {
3992 assert(
I.getNumOperands() == 3);
3993 assert(
I.getOperand(2).isReg());
3995 Register InputRegister =
I.getOperand(2).getReg();
3999 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
4002 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
4008 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
4009 .
addUse(
I.getOperand(2).getReg())
4014bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
4015 SPIRVTypeInst ResType,
4018 assert(
I.getNumOperands() == 3);
4019 assert(
I.getOperand(2).isReg());
4021 Register InputRegister =
I.getOperand(2).getReg();
4027 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
4038bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
4039 SPIRVTypeInst ResType,
4046 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
4051 : SPIRV::OpUConvert;
4053 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4056 ShiftOp = SPIRV::OpShiftRightLogicalV;
4061 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4062 TII.get(SPIRV::OpConstantComposite))
4065 for (
unsigned It = 0; It <
N; ++It)
4069 ShiftConst = CompositeReg;
4074 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
4079 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
4084 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
4089 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
4092bool SPIRVInstructionSelector::handle64BitOverflow(
4094 unsigned int Opcode,
4101 "handle64BitOverflow should only be used for integer types");
4103 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4105 MachineIRBuilder MIRBuilder(
I);
4107 SPIRVTypeInst I64x2Type =
4109 SPIRVTypeInst Vec2ResType =
4112 std::vector<Register> PartialRegs;
4114 unsigned CurrentComponent = 0;
4115 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4119 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4120 TII.get(SPIRV::OpVectorShuffle))
4125 .
addImm(CurrentComponent)
4126 .
addImm(CurrentComponent + 1);
4136 PartialRegs.push_back(SubVecReg);
4139 if (CurrentComponent != ComponentCount) {
4145 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4146 SPIRV::OpVectorExtractDynamic))
4155 PartialRegs.push_back(FinalElemResReg);
4159 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4160 SPIRV::OpCompositeConstruct);
4163bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4164 SPIRVTypeInst ResType,
4168 if (ComponentCount > 2)
4169 return handle64BitOverflow(
4170 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4172 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4174 MachineIRBuilder MIRBuilder(
I);
4178 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4182 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4187 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4194 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4195 TII.get(SPIRV::OpVectorShuffle))
4200 for (
unsigned J = 0; J < ComponentCount; ++J) {
4207 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4210bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4211 SPIRVTypeInst ResType,
4215 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4223bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4224 SPIRVTypeInst ResType,
4225 MachineInstr &
I)
const {
4226 Register OpReg =
I.getOperand(1).getReg();
4235 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4237 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4239 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4241 return SPIRVInstructionSelector::diagnoseUnsupported(
4242 I,
"G_BITREVERSE only support 16,32,64 bits.");
4246 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4257 unsigned AndOp = SPIRV::OpBitwiseAndS;
4258 unsigned OrOp = SPIRV::OpBitwiseOrS;
4259 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4260 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4261 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4263 AndOp = SPIRV::OpBitwiseAndV;
4264 OrOp = SPIRV::OpBitwiseOrV;
4265 ShlOp = SPIRV::OpShiftLeftLogicalV;
4266 ShrOp = SPIRV::OpShiftRightLogicalV;
4272 const unsigned Shift) ->
Register {
4275 (ResType->
getOpcode() != SPIRV::OpTypeVectorIdEXT ||
4282 Register MaskReg = CreateConst(Mask);
4283 Register ShiftReg = CreateConst(Shift);
4290 if (!selectOpWithSrcs(
T1, ResType,
I, {
Input, ShiftReg}, ShrOp) ||
4291 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4292 !selectOpWithSrcs(T3, ResType,
I, {
Input, MaskReg}, AndOp) ||
4293 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4294 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4303 while ((Shift >>= 1) > 0) {
4310 return BuildCOPY(ResVReg, Result,
I);
4313bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4314 SPIRVTypeInst ResType,
4315 MachineInstr &
I)
const {
4316 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4317 "G_FREEZE must define and use a register");
4318 Register OpReg =
I.getOperand(1).getReg();
4322 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4336 if (
Def->getOpcode() == TargetOpcode::COPY)
4339 switch (
Def->getOpcode()) {
4340 case SPIRV::ASSIGN_TYPE:
4341 if (MachineInstr *AssignToDef =
4343 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4344 Reg =
Def->getOperand(2).getReg();
4347 case SPIRV::OpUndef:
4348 Reg =
Def->getOperand(1).getReg();
4351 unsigned DestOpCode;
4353 DestOpCode = SPIRV::OpConstantNull;
4354 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4355 "static undef/poison lowered to OpConstantNull\n");
4357 DestOpCode = TargetOpcode::COPY;
4359 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4360 "skipped, lowered as a copy of the operand\n");
4362 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4363 .
addDef(
I.getOperand(0).getReg())
4371bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4372 SPIRVTypeInst ResType,
4373 MachineInstr &
I)
const {
4377 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4381 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4386 for (
unsigned i =
I.getNumExplicitDefs();
4387 i <
I.getNumExplicitOperands() && IsConst; ++i)
4392 return diagnoseUnsupported(
4393 I,
"There must be at least two constituent operands in a vector");
4398 for (
unsigned i =
I.getNumExplicitDefs();
4399 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4405 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4412 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4413 TII.get(IsConst ? SPIRV::OpConstantComposite
4414 : SPIRV::OpCompositeConstruct))
4417 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4418 MIB.
addUse(
I.getOperand(i).getReg());
4423bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4424 SPIRVTypeInst ResType,
4425 MachineInstr &
I)
const {
4429 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4434 unsigned OpIdx =
I.getNumExplicitDefs();
4435 if (!
I.getOperand(OpIdx).isReg())
4439 Register OpReg =
I.getOperand(OpIdx).getReg();
4443 return diagnoseUnsupported(
4444 I,
"There must be at least two constituent operands in a vector");
4447 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4448 TII.get(IsConst ? SPIRV::OpConstantComposite
4449 : SPIRV::OpCompositeConstruct))
4452 for (
unsigned i = 0; i <
N; ++i)
4458bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4459 SPIRVTypeInst ResType,
4460 MachineInstr &
I)
const {
4466 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4468 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4469 TII.get(SPIRV::OpCompositeConstruct))
4472 for (
unsigned OpIdx =
I.getNumExplicitDefs();
4473 OpIdx <
I.getNumExplicitOperands(); ++OpIdx)
4474 MIB.
addUse(
I.getOperand(OpIdx).getReg());
4479bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4480 SPIRVTypeInst ResType,
4481 MachineInstr &
I)
const {
4487 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4489 Opcode = SPIRV::OpDemoteToHelperInvocation;
4491 Opcode = SPIRV::OpKill;
4496 ToErase.eraseFromParent();
4505bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4506 SPIRVTypeInst ResType,
unsigned CmpOpc,
4507 MachineInstr &
I)
const {
4508 Register Cmp0 =
I.getOperand(2).getReg();
4509 Register Cmp1 =
I.getOperand(3).getReg();
4512 "CMP operands should have the same type");
4513 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4523bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4524 SPIRVTypeInst ResType,
4525 MachineInstr &
I)
const {
4526 auto Pred =
I.getOperand(1).getPredicate();
4529 Register CmpOperand =
I.getOperand(2).getReg();
4531 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4536 Register Op1 =
I.getOperand(3).getReg();
4540 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4545 I.getOperand(3).setReg(NewOp1);
4551 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4556Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4557 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4559 SPIRVTypeInst SpvI32Ty =
4561 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4566 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4567 MachineInstr *
MI =
nullptr;
4571 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4575 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4576 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4582 GR.
add(ConstInt,
MI);
4587bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4588 SPIRVTypeInst ResType,
4589 MachineInstr &
I)
const {
4590 auto Pred =
I.getOperand(1).getPredicate();
4598 for (
unsigned Idx = 0;
Idx < 2; ++
Idx) {
4600 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
4603 .
addUse(
I.getOperand(2 + Idx).getReg())
4610 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLogicalOr))
4617 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLogicalNot))
4625 return selectCmp(ResVReg, ResType, CmpOp,
I);
4628bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4629 SPIRVTypeInst ResType,
4630 MachineInstr &
I)
const {
4632 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4642 MachineIRBuilder MIRBuilder(
I);
4649 APFloat ConstVal(3.3219280948873623);
4653 APFloat::rmNearestTiesToEven, &LosesInfo);
4658 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
4660 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4661 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4663 if (!selectExtInst(ResVReg, ResType,
I,
4664 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4674Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4675 MachineInstr &
I)
const {
4683bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4689 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4697 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4700 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4701 Def->getOpcode() == SPIRV::OpConstantI)
4714 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4715 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4717 Intrinsic::spv_const_composite)) {
4718 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4719 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4720 if (!IsZero(
Def->getOperand(i).getReg()))
4729Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4730 MachineInstr &
I)
const {
4739Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4740 MachineInstr &
I)
const {
4750 SPIRVTypeInst ResType,
4751 MachineInstr &
I)
const {
4762 MachineInstr &
I)
const {
4767 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4768 TII.get(SPIRV::OpCompositeConstruct))
4771 for (
unsigned J = 0; J < NumElts; ++J)
4777bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4778 SPIRVTypeInst ResType,
4779 MachineInstr &
I)
const {
4780 Register SelectFirstArg =
I.getOperand(2).getReg();
4781 Register SelectSecondArg =
I.getOperand(3).getReg();
4790 Register CondReg =
I.getOperand(1).getReg();
4791 bool IsScalarBool = GR.
isScalarOfType(CondReg, SPIRV::OpTypeBool);
4799 CondReg = buildVectorSplat(CondReg, NumElts,
I);
4800 IsScalarBool =
false;
4803 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4804 }
else if (IsPtrTy) {
4805 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4807 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4810 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4811 "boolean condition");
4813 Opcode = SPIRV::OpSelectSFSCond;
4814 }
else if (IsPtrTy) {
4815 Opcode = SPIRV::OpSelectSPSCond;
4817 Opcode = SPIRV::OpSelectSISCond;
4820 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4832bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4833 SPIRVTypeInst ResType,
4835 MachineInstr &InsertAt,
4836 bool IsSigned)
const {
4838 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4839 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4840 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4842 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4854bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4855 SPIRVTypeInst ResType,
4856 MachineInstr &
I,
bool IsSigned,
4857 unsigned Opcode)
const {
4858 Register SrcReg =
I.getOperand(1).getReg();
4869 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4871 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4874bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4875 SPIRVTypeInst ResType, MachineInstr &
I,
4876 bool IsSigned)
const {
4877 Register SrcReg =
I.getOperand(1).getReg();
4879 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4883 if (ResType == SrcType)
4884 return BuildCOPY(ResVReg, SrcReg,
I);
4886 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4887 return selectUnOp(ResVReg, ResType,
I, Opcode);
4890bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4891 SPIRVTypeInst ResType,
4893 bool IsSigned)
const {
4894 MachineIRBuilder MIRBuilder(
I);
4895 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
4900 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4908 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4911 .
addUse(
I.getOperand(1).getReg())
4912 .
addUse(
I.getOperand(2).getReg())
4917 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4920 .
addUse(
I.getOperand(1).getReg())
4921 .
addUse(
I.getOperand(2).getReg())
4929 unsigned SelectOpcode =
4930 (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4932 ? SPIRV::OpSelectVIVCond
4933 : SPIRV::OpSelectSISCond;
4938 .
addUse(buildOnesVal(
true, ResType,
I))
4939 .
addUse(buildZerosVal(ResType,
I))
4946 .
addUse(buildOnesVal(
false, ResType,
I))
4951bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4954 SPIRVTypeInst IntTy,
4955 SPIRVTypeInst BoolTy)
const {
4959 isVectorType(IntTy) ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4961 Register One = buildOnesVal(
false, IntTy,
I);
4969 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4978bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4979 SPIRVTypeInst ResType,
4980 MachineInstr &
I)
const {
4981 Register IntReg =
I.getOperand(1).getReg();
4984 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4985 if (ArgType == ResType)
4986 return BuildCOPY(ResVReg, IntReg,
I);
4988 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4989 return selectUnOp(ResVReg, ResType,
I, Opcode);
4992bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4993 SPIRVTypeInst ResType,
4994 MachineInstr &
I)
const {
4995 unsigned Opcode =
I.getOpcode();
4996 unsigned TpOpcode = ResType->
getOpcode();
4998 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
4999 assert(Opcode == TargetOpcode::G_CONSTANT &&
5000 I.getOperand(1).getCImm()->isZero());
5001 MachineBasicBlock &DepMBB =
I.getMF()->front();
5004 }
else if (TpOpcode == SPIRV::OpTypeVectorIdEXT) {
5009 "Expected <1 x T> Vector!");
5010 if (Opcode == TargetOpcode::G_FCONSTANT)
5016 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
5024 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
5027bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
5028 SPIRVTypeInst ResType,
5029 MachineInstr &
I)
const {
5030 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5037bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
5038 SPIRVTypeInst ResType,
5039 MachineInstr &
I)
const {
5041 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
5045 .
addUse(
I.getOperand(3).getReg())
5047 .
addUse(
I.getOperand(2).getReg());
5048 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
5054bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
5055 SPIRVTypeInst ResType,
5056 MachineInstr &
I)
const {
5057 Type *MaybeResTy =
nullptr;
5062 "Expected aggregate type for extractv instruction");
5064 SPIRV::AccessQualifier::ReadWrite,
false);
5068 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5071 .
addUse(
I.getOperand(2).getReg());
5072 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
5078bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
5079 SPIRVTypeInst ResType,
5080 MachineInstr &
I)
const {
5082 return selectInsertVal(ResVReg, ResType,
I);
5084 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
5087 .
addUse(
I.getOperand(2).getReg())
5088 .
addUse(
I.getOperand(3).getReg())
5089 .
addUse(
I.getOperand(4).getReg())
5094bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
5095 SPIRVTypeInst ResType,
5096 MachineInstr &
I)
const {
5098 return selectExtractVal(ResVReg, ResType,
I);
5100 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
5103 .
addUse(
I.getOperand(2).getReg())
5104 .
addUse(
I.getOperand(3).getReg())
5109bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
5110 SPIRVTypeInst ResType,
5111 MachineInstr &
I)
const {
5112 const bool IsGEPInBounds =
I.getOperand(2).getImm();
5115 const bool UseUntypedPointers =
5116 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
5121 if (UseUntypedPointers) {
5123 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
5124 : SPIRV::OpUntypedAccessChainKHR;
5126 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
5127 : SPIRV::OpUntypedPtrAccessChainKHR;
5136 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
5138 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
5139 : SPIRV::OpPtrAccessChain;
5144 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5149 if (UseUntypedPointers) {
5164 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5165 Def->getOperand(1).isReg())
5167 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5168 if (
const auto *GVar =
5171 SPIRV::AccessQualifier::ReadWrite,
5175 return diagnoseUnsupported(
5176 I,
"could not deduce the base type of an untyped access chain");
5181 Res.addUse(BaseReg);
5183 const bool IsAccessChainOpcode =
5184 (Opcode == SPIRV::OpAccessChain ||
5185 Opcode == SPIRV::OpInBoundsAccessChain ||
5186 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5187 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5191 "Cannot translate GEP to OpAccessChain.");
5194 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5195 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5196 Res.addUse(
I.getOperand(i).getReg());
5197 Res.constrainAllUses(
TII,
TRI, RBI);
5206 if (Extract.
getOpcode() == SPIRV::OpCompositeExtract) {
5212 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5218 TII.get(SPIRV::OpCompositeInsert))
5232bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5234 unsigned Lim =
I.getNumExplicitOperands();
5235 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5236 Register OpReg =
I.getOperand(i).getReg();
5240 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5241 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5254 SPIRVTypeInst WrapType = OpType;
5255 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5257 SPIRV::StorageClass::CodeSectionINTEL) {
5259 SPIRV::StorageClass::Function,
I);
5266 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5267 TII.get(SPIRV::OpSpecConstantOp))
5270 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5272 GR.
add(OpDefine, MIB);
5278bool SPIRVInstructionSelector::selectDerivativeInst(
5279 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5280 const unsigned DPdOpCode)
const {
5283 if (!errorIfInstrOutsideShader(
I))
5289 Register SrcReg =
I.getOperand(2).getReg();
5294 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5297 .
addUse(
I.getOperand(2).getReg());
5299 MachineIRBuilder MIRBuilder(
I);
5302 if (componentCount != 1)
5310 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5315 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5320 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5328bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5329 SPIRVTypeInst ResType,
5330 MachineInstr &
I)
const {
5334 case Intrinsic::spv_load:
5335 return selectLoad(ResVReg, ResType,
I);
5336 case Intrinsic::spv_atomic_load:
5337 return selectAtomicLoad(ResVReg, ResType,
I);
5338 case Intrinsic::spv_store:
5339 return selectStore(
I);
5340 case Intrinsic::spv_atomic_store:
5341 return selectAtomicStore(
I);
5342 case Intrinsic::spv_extractv:
5343 return selectExtractVal(ResVReg, ResType,
I);
5344 case Intrinsic::spv_insertv:
5345 return selectInsertVal(ResVReg, ResType,
I);
5346 case Intrinsic::spv_extractelt:
5347 return selectExtractElt(ResVReg, ResType,
I);
5348 case Intrinsic::spv_insertelt:
5349 return selectInsertElt(ResVReg, ResType,
I);
5350 case Intrinsic::spv_gep:
5351 return selectGEP(ResVReg, ResType,
I);
5352 case Intrinsic::spv_bitcast: {
5353 Register OpReg =
I.getOperand(2).getReg();
5354 SPIRVTypeInst OpType =
5358 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5360 case Intrinsic::spv_unref_global:
5361 case Intrinsic::spv_init_global: {
5367 Register GVarVReg =
MI->getOperand(0).getReg();
5368 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5373 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5375 MI->eraseFromParent();
5379 case Intrinsic::spv_undef: {
5380 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5386 case Intrinsic::spv_poison:
5387 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5392 case Intrinsic::spv_freeze:
5393 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5396 .
addUse(
I.getOperand(2).getReg())
5399 case Intrinsic::spv_named_boolean_spec_constant: {
5400 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5401 : SPIRV::OpSpecConstantFalse;
5403 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5404 .
addDef(
I.getOperand(0).getReg())
5407 unsigned SpecId =
I.getOperand(2).getImm();
5409 SPIRV::Decoration::SpecId, {SpecId});
5413 case Intrinsic::spv_const_composite: {
5415 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5421 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5423 std::function<bool(
Register)> HasSpecConstOperand =
5433 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5434 J < Def->getNumExplicitOperands(); ++J) {
5435 if (
Def->getOperand(J).isReg() &&
5436 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5442 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5443 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5444 : SPIRV::OpConstantComposite;
5445 unsigned ContinuedOpc = HasSpecConst
5446 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5447 : SPIRV::OpConstantCompositeContinuedINTEL;
5448 MachineIRBuilder MIR(
I);
5450 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5452 for (
auto *Instr : Instructions) {
5453 Instr->setDebugLoc(
I.getDebugLoc());
5458 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5465 case Intrinsic::spv_assign_name: {
5466 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5467 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5468 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5469 i <
I.getNumExplicitOperands(); ++i) {
5470 MIB.
addImm(
I.getOperand(i).getImm());
5475 case Intrinsic::spv_switch: {
5476 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5477 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5478 if (
I.getOperand(i).isReg())
5479 MIB.
addReg(
I.getOperand(i).getReg());
5480 else if (
I.getOperand(i).isCImm())
5481 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5482 else if (
I.getOperand(i).isMBB())
5483 MIB.
addMBB(
I.getOperand(i).getMBB());
5490 case Intrinsic::spv_loop_merge: {
5491 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5492 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5493 if (
I.getOperand(i).isMBB())
5494 MIB.
addMBB(
I.getOperand(i).getMBB());
5501 case Intrinsic::spv_loop_control_intel: {
5503 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5504 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5509 case Intrinsic::spv_selection_merge: {
5511 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5512 assert(
I.getOperand(1).isMBB() &&
5513 "operand 1 to spv_selection_merge must be a basic block");
5514 MIB.
addMBB(
I.getOperand(1).getMBB());
5515 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5519 case Intrinsic::spv_cmpxchg:
5520 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5521 case Intrinsic::spv_unreachable:
5522 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5525 case Intrinsic::spv_abort:
5526 return selectAbort(
I);
5527 case Intrinsic::spv_alloca:
5528 return selectFrameIndex(ResVReg, ResType,
I);
5529 case Intrinsic::spv_alloca_array:
5530 return selectAllocaArray(ResVReg, ResType,
I);
5531 case Intrinsic::spv_assume:
5533 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5534 .
addUse(
I.getOperand(1).getReg())
5539 case Intrinsic::spv_expect:
5541 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5544 .
addUse(
I.getOperand(2).getReg())
5545 .
addUse(
I.getOperand(3).getReg())
5550 case Intrinsic::arithmetic_fence:
5551 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5552 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5555 .
addUse(
I.getOperand(2).getReg())
5559 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5561 case Intrinsic::spv_thread_id:
5567 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5569 case Intrinsic::spv_thread_id_in_group:
5575 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5577 case Intrinsic::spv_group_id:
5583 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5585 case Intrinsic::spv_flattened_thread_id_in_group:
5592 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5594 case Intrinsic::spv_workgroup_size:
5595 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5597 case Intrinsic::spv_global_size:
5598 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5600 case Intrinsic::spv_global_offset:
5601 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5603 case Intrinsic::spv_num_workgroups:
5604 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5606 case Intrinsic::spv_subgroup_size:
5607 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5609 case Intrinsic::spv_num_subgroups:
5610 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5612 case Intrinsic::spv_subgroup_id:
5613 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5614 case Intrinsic::spv_subgroup_local_invocation_id:
5615 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5616 ResVReg, ResType,
I);
5617 case Intrinsic::spv_subgroup_max_size:
5618 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5620 case Intrinsic::spv_fdot:
5621 return selectFloatDot(ResVReg, ResType,
I);
5622 case Intrinsic::spv_udot:
5623 case Intrinsic::spv_sdot:
5624 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5626 return selectIntegerDot(ResVReg, ResType,
I,
5627 IID == Intrinsic::spv_sdot);
5628 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5629 case Intrinsic::spv_dot4add_i8packed:
5630 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5632 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5633 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5634 case Intrinsic::spv_dot4add_u8packed:
5635 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5637 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5638 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5639 case Intrinsic::spv_all:
5640 return selectAll(ResVReg, ResType,
I);
5641 case Intrinsic::spv_any:
5642 return selectAny(ResVReg, ResType,
I);
5643 case Intrinsic::spv_distance:
5644 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5645 case Intrinsic::spv_lerp:
5646 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5647 case Intrinsic::spv_length:
5648 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5649 case Intrinsic::spv_degrees:
5650 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5651 case Intrinsic::spv_faceforward:
5652 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5653 case Intrinsic::spv_frac:
5654 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5655 case Intrinsic::spv_isinf:
5656 return selectOpIsInf(ResVReg, ResType,
I);
5657 case Intrinsic::spv_isnan:
5658 return selectOpIsNan(ResVReg, ResType,
I);
5659 case Intrinsic::spv_isfinite:
5660 return selectOpIsFinite(ResVReg, ResType,
I);
5661 case Intrinsic::spv_isnormal:
5662 return selectOpIsNormal(ResVReg, ResType,
I);
5663 case Intrinsic::spv_normalize:
5664 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5665 case Intrinsic::spv_refract:
5666 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5667 case Intrinsic::spv_reflect:
5668 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5669 case Intrinsic::spv_rsqrt:
5670 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5671 case Intrinsic::spv_sign:
5672 return selectSign(ResVReg, ResType,
I);
5673 case Intrinsic::spv_smoothstep:
5674 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5675 case Intrinsic::spv_firstbituhigh:
5676 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5677 case Intrinsic::spv_firstbitshigh:
5678 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5679 case Intrinsic::spv_firstbitlow:
5680 return selectFirstBitLow(ResVReg, ResType,
I);
5681 case Intrinsic::spv_all_memory_barrier:
5682 return selectBarrierInst(
I, SPIRV::Scope::Device,
5683 SPIRV::MemorySemantics::UniformMemory |
5684 SPIRV::MemorySemantics::ImageMemory |
5685 SPIRV::MemorySemantics::WorkgroupMemory,
5687 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5688 return selectBarrierInst(
I, SPIRV::Scope::Device,
5689 SPIRV::MemorySemantics::UniformMemory |
5690 SPIRV::MemorySemantics::ImageMemory |
5691 SPIRV::MemorySemantics::WorkgroupMemory,
5693 case Intrinsic::spv_device_memory_barrier:
5694 return selectBarrierInst(
I, SPIRV::Scope::Device,
5695 SPIRV::MemorySemantics::UniformMemory |
5696 SPIRV::MemorySemantics::ImageMemory,
5698 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5699 return selectBarrierInst(
I, SPIRV::Scope::Device,
5700 SPIRV::MemorySemantics::UniformMemory |
5701 SPIRV::MemorySemantics::ImageMemory,
5703 case Intrinsic::spv_group_memory_barrier:
5704 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5705 SPIRV::MemorySemantics::WorkgroupMemory,
5707 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5708 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5709 SPIRV::MemorySemantics::WorkgroupMemory,
5711 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5712 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5713 SPIRV::StorageClass::StorageClass ResSC =
5716 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5717 "from the Generic storage class");
5718 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5726 case Intrinsic::spv_lifetime_start:
5727 case Intrinsic::spv_lifetime_end: {
5728 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5729 : SPIRV::OpLifetimeStop;
5730 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5731 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5740 case Intrinsic::spv_saturate:
5741 return selectSaturate(ResVReg, ResType,
I);
5742 case Intrinsic::spv_nclamp:
5743 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5744 case Intrinsic::spv_uclamp:
5745 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5746 case Intrinsic::spv_sclamp:
5747 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5748 case Intrinsic::spv_subgroup_prefix_bit_count:
5749 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5750 case Intrinsic::spv_wave_active_countbits:
5751 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5752 case Intrinsic::spv_wave_all_equal:
5753 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5754 case Intrinsic::spv_wave_all:
5755 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5756 case Intrinsic::spv_wave_any:
5757 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5758 case Intrinsic::spv_subgroup_ballot:
5759 return selectWaveOpInst(ResVReg, ResType,
I,
5760 SPIRV::OpGroupNonUniformBallot);
5761 case Intrinsic::spv_wave_is_first_lane:
5762 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5763 case Intrinsic::spv_wave_reduce_or:
5764 return selectWaveReduceOp(ResVReg, ResType,
I,
5765 SPIRV::OpGroupNonUniformBitwiseOr);
5766 case Intrinsic::spv_wave_reduce_xor:
5767 return selectWaveReduceOp(ResVReg, ResType,
I,
5768 SPIRV::OpGroupNonUniformBitwiseXor);
5769 case Intrinsic::spv_wave_reduce_and:
5770 return selectWaveReduceOp(ResVReg, ResType,
I,
5771 SPIRV::OpGroupNonUniformBitwiseAnd);
5772 case Intrinsic::spv_wave_reduce_umax:
5773 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5774 case Intrinsic::spv_wave_reduce_max:
5775 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5776 case Intrinsic::spv_wave_reduce_umin:
5777 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5778 case Intrinsic::spv_wave_reduce_min:
5779 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5780 case Intrinsic::spv_wave_reduce_sum:
5781 return selectWaveReduceSum(ResVReg, ResType,
I);
5782 case Intrinsic::spv_wave_product:
5783 return selectWaveReduceProduct(ResVReg, ResType,
I);
5784 case Intrinsic::spv_wave_readlane:
5785 return selectWaveOpInst(ResVReg, ResType,
I,
5786 SPIRV::OpGroupNonUniformShuffle);
5787 case Intrinsic::spv_wave_readlane_first:
5788 return selectWaveOpInst(ResVReg, ResType,
I,
5789 SPIRV::OpGroupNonUniformBroadcastFirst);
5790 case Intrinsic::spv_wave_prefix_sum:
5791 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5792 case Intrinsic::spv_wave_prefix_product:
5793 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5794 case Intrinsic::spv_quad_read_across_x: {
5795 return selectQuadSwap(ResVReg, ResType,
I, 0);
5797 case Intrinsic::spv_quad_read_across_y: {
5798 return selectQuadSwap(ResVReg, ResType,
I, 1);
5800 case Intrinsic::spv_quad_read_across_diagonal: {
5801 return selectQuadSwap(ResVReg, ResType,
I, 2);
5803 case Intrinsic::spv_radians:
5804 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5808 case Intrinsic::instrprof_increment:
5809 case Intrinsic::instrprof_increment_step:
5810 case Intrinsic::instrprof_value_profile:
5813 case Intrinsic::spv_value_md:
5815 case Intrinsic::spv_resource_handlefrombinding: {
5816 return selectHandleFromBinding(ResVReg, ResType,
I);
5818 case Intrinsic::spv_resource_counterhandlefrombinding:
5819 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5820 case Intrinsic::spv_resource_updatecounter:
5821 return selectUpdateCounter(ResVReg, ResType,
I);
5822 case Intrinsic::spv_resource_store_typedbuffer: {
5823 return selectImageWriteIntrinsic(
I);
5825 case Intrinsic::spv_resource_load_typedbuffer: {
5826 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5828 case Intrinsic::spv_resource_load_level: {
5829 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5831 case Intrinsic::spv_resource_getdimensions_x:
5832 case Intrinsic::spv_resource_getdimensions_xy:
5833 case Intrinsic::spv_resource_getdimensions_xyz: {
5834 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5836 case Intrinsic::spv_resource_getdimensions_levels_x:
5837 case Intrinsic::spv_resource_getdimensions_levels_xy:
5838 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5839 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5841 case Intrinsic::spv_resource_getdimensions_ms_xy:
5842 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5843 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5845 case Intrinsic::spv_resource_calculate_lod:
5846 case Intrinsic::spv_resource_calculate_lod_unclamped:
5847 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5848 case Intrinsic::spv_resource_sample:
5849 case Intrinsic::spv_resource_sample_clamp:
5850 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5851 case Intrinsic::spv_resource_samplebias:
5852 case Intrinsic::spv_resource_samplebias_clamp:
5853 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5854 case Intrinsic::spv_resource_samplegrad:
5855 case Intrinsic::spv_resource_samplegrad_clamp:
5856 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5857 case Intrinsic::spv_resource_samplelevel:
5858 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5859 case Intrinsic::spv_resource_samplecmp:
5860 case Intrinsic::spv_resource_samplecmp_clamp:
5861 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5862 case Intrinsic::spv_resource_samplecmplevelzero:
5863 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5864 case Intrinsic::spv_resource_gather:
5865 case Intrinsic::spv_resource_gather_cmp:
5866 return selectGatherIntrinsic(ResVReg, ResType,
I);
5867 case Intrinsic::spv_resource_getbasepointer:
5868 case Intrinsic::spv_resource_getpointer: {
5869 return selectResourceGetPointer(ResVReg, ResType,
I);
5871 case Intrinsic::spv_pushconstant_getpointer: {
5872 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5874 case Intrinsic::spv_discard: {
5875 return selectDiscard(ResVReg, ResType,
I);
5877 case Intrinsic::spv_resource_nonuniformindex: {
5878 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5880 case Intrinsic::spv_unpackhalf2x16: {
5881 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5883 case Intrinsic::spv_packhalf2x16: {
5884 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5886 case Intrinsic::spv_ddx:
5887 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5888 case Intrinsic::spv_ddy:
5889 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5890 case Intrinsic::spv_ddx_coarse:
5891 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5892 case Intrinsic::spv_ddy_coarse:
5893 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5894 case Intrinsic::spv_ddx_fine:
5895 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5896 case Intrinsic::spv_ddy_fine:
5897 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5898 case Intrinsic::spv_fwidth:
5899 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5900 case Intrinsic::spv_masked_gather:
5901 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5902 return selectMaskedGather(ResVReg, ResType,
I);
5903 return diagnoseUnsupported(
5904 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5905 case Intrinsic::spv_masked_scatter:
5906 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5907 return selectMaskedScatter(
I);
5908 return diagnoseUnsupported(
5909 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5910 case Intrinsic::returnaddress:
5911 case Intrinsic::frameaddress: {
5913 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5920 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5925bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5926 SPIRVTypeInst ResType,
5927 MachineInstr &
I)
const {
5930 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5937bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5938 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5940 assert(Intr.getIntrinsicID() ==
5941 Intrinsic::spv_resource_counterhandlefrombinding);
5944 Register MainHandleReg = Intr.getOperand(2).getReg();
5946 assert(MainHandleDef->getIntrinsicID() ==
5947 Intrinsic::spv_resource_handlefrombinding);
5951 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(),
MRI);
5952 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5953 std::string CounterName =
5958 MachineIRBuilder MIRBuilder(
I);
5960 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5962 ArraySize, IndexReg, CounterName, MIRBuilder);
5964 return BuildCOPY(ResVReg, CounterVarReg,
I);
5967bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5968 SPIRVTypeInst ResType,
5969 MachineInstr &
I)
const {
5971 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5973 Register CounterHandleReg = Intr.getOperand(2).getReg();
5974 int64_t IncrVal = Intr.getOperand(3).getImm();
5981 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5982 assert(CounterVarPointeeType &&
5983 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5984 "Counter variable must be a struct");
5986 SPIRV::StorageClass::StorageBuffer &&
5987 "Counter variable must be in the storage buffer storage class");
5989 "Counter variable must have exactly 1 member in the struct");
5990 const SPIRVTypeInst MemberType =
5993 "Counter variable struct must have a single i32 member");
5997 MachineIRBuilder MIRBuilder(
I);
5999 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
6002 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
6008 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6011 .
addUse(CounterHandleReg)
6017 Register Scope = buildI32ConstantInEntryBlock(SPIRV::Scope::Device,
I);
6019 buildI32ConstantInEntryBlock(SPIRV::MemorySemantics::None,
I);
6022 buildI32ConstantInEntryBlock(
static_cast<uint32_t
>(IncrVal),
I);
6025 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
6034 return BuildCOPY(ResVReg, AtomicRes,
I);
6042 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
6050bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
6051 SPIRVTypeInst ResType,
6052 MachineInstr &
I)
const {
6060 Register ImageReg =
I.getOperand(2).getReg();
6068 Register IdxReg =
I.getOperand(3).getReg();
6070 MachineInstr &Pos =
I;
6072 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
6076bool SPIRVInstructionSelector::generateSampleImage(
6079 DebugLoc Loc, MachineInstr &Pos)
const {
6090 if (!loadHandleBeforePosition(NewSamplerReg,
6096 MachineIRBuilder MIRBuilder(Pos);
6109 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
6110 ImOps.Lod.has_value();
6111 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
6112 : SPIRV::OpImageSampleImplicitLod;
6114 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
6115 : SPIRV::OpImageSampleDrefImplicitLod;
6124 MIB.
addUse(*ImOps.Compare);
6126 uint32_t ImageOperands = 0;
6128 ImageOperands |= SPIRV::ImageOperand::Bias;
6130 ImageOperands |= SPIRV::ImageOperand::Lod;
6131 if (ImOps.GradX && ImOps.GradY)
6132 ImageOperands |= SPIRV::ImageOperand::Grad;
6133 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
6135 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6138 "Non-constant offsets are not supported in sample instructions.");
6143 ImageOperands |= SPIRV::ImageOperand::MinLod;
6145 if (ImageOperands != 0) {
6146 MIB.
addImm(ImageOperands);
6147 if (ImageOperands & SPIRV::ImageOperand::Bias)
6149 if (ImageOperands & SPIRV::ImageOperand::Lod)
6151 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6152 MIB.
addUse(*ImOps.GradX);
6153 MIB.
addUse(*ImOps.GradY);
6156 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6157 MIB.
addUse(*ImOps.Offset);
6158 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6159 MIB.
addUse(*ImOps.MinLod);
6166bool SPIRVInstructionSelector::selectImageQuerySize(
6168 std::optional<Register> LodReg)
const {
6171 "ImageReg is not an image type.");
6173 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6175 unsigned NumComponents = 0;
6177 case SPIRV::Dim::DIM_1D:
6178 case SPIRV::Dim::DIM_Buffer:
6179 NumComponents =
IsArray ? 2 : 1;
6181 case SPIRV::Dim::DIM_2D:
6182 case SPIRV::Dim::DIM_Cube:
6183 case SPIRV::Dim::DIM_Rect:
6184 NumComponents =
IsArray ? 3 : 2;
6186 case SPIRV::Dim::DIM_3D:
6190 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6200 if (!LodReg && IsSampled && !IsMultisampled &&
6201 Dim != SPIRV::Dim::DIM_Buffer && Dim != SPIRV::Dim::DIM_Rect)
6204 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6205 SPIRVTypeInst ResType =
6210 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6220bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6221 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6222 Register ImageReg =
I.getOperand(2).getReg();
6229 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6232bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6233 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6234 Register ImageReg =
I.getOperand(2).getReg();
6243 Register LodReg =
I.getOperand(3).getReg();
6246 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6248 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6255 TII.get(SPIRV::OpImageQueryLevels))
6262 TII.get(SPIRV::OpCompositeConstruct))
6272bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6273 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6274 Register ImageReg =
I.getOperand(2).getReg();
6285 "OpImageQuerySamples requires a multisampled image");
6287 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6295 TII.get(SPIRV::OpImageQuerySamples))
6302 TII.get(SPIRV::OpCompositeConstruct))
6312bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6313 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6314 Register ImageReg =
I.getOperand(2).getReg();
6315 Register SamplerReg =
I.getOperand(3).getReg();
6316 Register CoordinateReg =
I.getOperand(4).getReg();
6332 if (!loadHandleBeforePosition(
6337 MachineIRBuilder MIRBuilder(
I);
6343 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6353 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6360 unsigned ExtractedIndex =
6362 Intrinsic::spv_resource_calculate_lod_unclamped
6366 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6367 TII.get(SPIRV::OpCompositeExtract))
6377bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6378 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6379 Register ImageReg =
I.getOperand(2).getReg();
6380 Register SamplerReg =
I.getOperand(3).getReg();
6381 Register CoordinateReg =
I.getOperand(4).getReg();
6382 ImageOperands ImOps;
6383 if (
I.getNumOperands() > 5)
6384 ImOps.Offset =
I.getOperand(5).getReg();
6385 if (
I.getNumOperands() > 6)
6386 ImOps.MinLod =
I.getOperand(6).getReg();
6387 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6388 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6391bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6392 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6393 Register ImageReg =
I.getOperand(2).getReg();
6394 Register SamplerReg =
I.getOperand(3).getReg();
6395 Register CoordinateReg =
I.getOperand(4).getReg();
6396 ImageOperands ImOps;
6397 ImOps.Bias =
I.getOperand(5).getReg();
6398 if (
I.getNumOperands() > 6)
6399 ImOps.Offset =
I.getOperand(6).getReg();
6400 if (
I.getNumOperands() > 7)
6401 ImOps.MinLod =
I.getOperand(7).getReg();
6402 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6403 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6406bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6407 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6408 Register ImageReg =
I.getOperand(2).getReg();
6409 Register SamplerReg =
I.getOperand(3).getReg();
6410 Register CoordinateReg =
I.getOperand(4).getReg();
6411 ImageOperands ImOps;
6412 ImOps.GradX =
I.getOperand(5).getReg();
6413 ImOps.GradY =
I.getOperand(6).getReg();
6414 if (
I.getNumOperands() > 7)
6415 ImOps.Offset =
I.getOperand(7).getReg();
6416 if (
I.getNumOperands() > 8)
6417 ImOps.MinLod =
I.getOperand(8).getReg();
6418 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6419 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6422bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6423 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6424 Register ImageReg =
I.getOperand(2).getReg();
6425 Register SamplerReg =
I.getOperand(3).getReg();
6426 Register CoordinateReg =
I.getOperand(4).getReg();
6427 ImageOperands ImOps;
6428 ImOps.Lod =
I.getOperand(5).getReg();
6429 if (
I.getNumOperands() > 6)
6430 ImOps.Offset =
I.getOperand(6).getReg();
6431 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6432 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6435bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6436 SPIRVTypeInst ResType,
6437 MachineInstr &
I)
const {
6438 Register ImageReg =
I.getOperand(2).getReg();
6439 Register SamplerReg =
I.getOperand(3).getReg();
6440 Register CoordinateReg =
I.getOperand(4).getReg();
6441 ImageOperands ImOps;
6442 ImOps.Compare =
I.getOperand(5).getReg();
6443 if (
I.getNumOperands() > 6)
6444 ImOps.Offset =
I.getOperand(6).getReg();
6445 if (
I.getNumOperands() > 7)
6446 ImOps.MinLod =
I.getOperand(7).getReg();
6447 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6448 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6451bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6452 SPIRVTypeInst ResType,
6453 MachineInstr &
I)
const {
6454 Register ImageReg =
I.getOperand(2).getReg();
6455 Register CoordinateReg =
I.getOperand(3).getReg();
6456 Register LodReg =
I.getOperand(4).getReg();
6458 ImageOperands ImOps;
6460 if (
I.getNumOperands() > 5)
6461 ImOps.Offset =
I.getOperand(5).getReg();
6473 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6474 I.getDebugLoc(),
I, &ImOps);
6477bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6478 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6479 Register ImageReg =
I.getOperand(2).getReg();
6480 Register SamplerReg =
I.getOperand(3).getReg();
6481 Register CoordinateReg =
I.getOperand(4).getReg();
6482 ImageOperands ImOps;
6483 ImOps.Compare =
I.getOperand(5).getReg();
6484 if (
I.getNumOperands() > 6)
6485 ImOps.Offset =
I.getOperand(6).getReg();
6488 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6489 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6492bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6493 SPIRVTypeInst ResType,
6494 MachineInstr &
I)
const {
6495 Register ImageReg =
I.getOperand(2).getReg();
6496 Register SamplerReg =
I.getOperand(3).getReg();
6497 Register CoordinateReg =
I.getOperand(4).getReg();
6500 "ImageReg is not an image type.");
6505 ComponentOrCompareReg =
I.getOperand(5).getReg();
6506 OffsetReg =
I.getOperand(6).getReg();
6509 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6513 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6514 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6515 Dim != SPIRV::Dim::DIM_Rect) {
6517 "Gather operations are only supported for 2D, Cube, and Rect images.");
6524 if (!loadHandleBeforePosition(
6529 MachineIRBuilder MIRBuilder(
I);
6530 SPIRVTypeInst SampledImageType =
6535 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6543 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6545 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6547 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6552 .
addUse(ComponentOrCompareReg);
6554 uint32_t ImageOperands = 0;
6555 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6556 if (Dim == SPIRV::Dim::DIM_Cube) {
6558 "Gather operations with offset are not supported for Cube images.");
6562 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6564 ImageOperands |= SPIRV::ImageOperand::Offset;
6568 if (ImageOperands != 0) {
6569 MIB.
addImm(ImageOperands);
6571 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6579bool SPIRVInstructionSelector::generateImageReadOrFetch(
6582 const ImageOperands *ImOps)
const {
6585 "ImageReg is not an image type.");
6587 bool IsSignedInteger =
6592 bool IsFetch = (SampledOp.getImm() == 1);
6594 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6595 uint32_t ImageOperandsMask = 0;
6596 if (IsSignedInteger)
6597 ImageOperandsMask |= 0x1000;
6599 if (IsFetch && ImOps) {
6601 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6602 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6604 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6606 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6610 if (ImageOperandsMask != 0) {
6611 MIB.
addImm(ImageOperandsMask);
6612 if (IsFetch && ImOps) {
6615 if (ImOps->Offset &&
6616 (ImageOperandsMask &
6617 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6618 MIB.
addUse(*ImOps->Offset);
6627 SPIRVTypeInst SampledType =
6630 SPIRVTypeInst ReadType =
6631 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6632 bool ReadTypeMatchesResult = ReadType == ResType;
6634 Register ReadReg = ReadTypeMatchesResult
6640 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6646 BMI.constrainAllUses(
TII,
TRI, RBI);
6648 if (ReadTypeMatchesResult)
6661 if (ResultSize == 1) {
6670 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6673bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6674 SPIRVTypeInst ResType,
6675 MachineInstr &
I)
const {
6676 Register ResourcePtr =
I.getOperand(2).getReg();
6678 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6687 MachineIRBuilder MIRBuilder(
I);
6692 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6698 if (
I.getNumExplicitOperands() > 3) {
6699 Register IndexReg =
I.getOperand(3).getReg();
6706bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6707 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6712bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6713 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6714 Register ObjReg =
I.getOperand(2).getReg();
6715 if (!BuildCOPY(ResVReg, ObjReg,
I))
6725 decorateUsesAsNonUniform(ResVReg);
6729void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6732 {NonUniformReg,
nullptr}};
6733 llvm::SmallSet<Register, 8> Visited;
6734 while (WorkList.
size() > 0) {
6737 if (!Visited.
insert(CurrentReg).second)
6740 bool IsDecorated =
false;
6742 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6743 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6749 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6751 if (ResultReg == CurrentReg)
6759 MachineInstr &InsertPt =
6762 SPIRV::Decoration::NonUniformEXT, {});
6767bool SPIRVInstructionSelector::extractSubvector(
6769 MachineInstr &InsertionPoint)
const {
6771 [[maybe_unused]]
uint64_t InputSize =
6774 [[maybe_unused]]
bool IsLongVectorEXT =
6776 assert((InputSize > 1 || IsLongVectorEXT) &&
"The input must be a vector.");
6777 assert((ResultSize > 1 || IsLongVectorEXT) &&
"The result must be a vector.");
6778 assert(ResultSize < InputSize &&
6779 "Cannot extract more element than there are in the input.");
6786 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6795 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6797 TII.get(SPIRV::OpCompositeConstruct))
6801 for (
Register ComponentReg : ComponentRegisters)
6802 MIB.
addUse(ComponentReg);
6807bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6808 MachineInstr &
I)
const {
6815 Register ImageReg =
I.getOperand(1).getReg();
6823 Register CoordinateReg =
I.getOperand(2).getReg();
6824 Register DataReg =
I.getOperand(3).getReg();
6827 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6835Register SPIRVInstructionSelector::buildPointerToResource(
6836 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6837 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6838 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6840 if (ArraySize == 1) {
6841 SPIRVTypeInst PtrType =
6844 "SpirvResType did not have an explicit layout.");
6849 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6850 SPIRVTypeInst VarPointerType =
6853 VarPointerType, Set,
Binding, Name, MIRBuilder);
6855 SPIRVTypeInst ResPointerType =
6868bool SPIRVInstructionSelector::selectFirstBitSet16(
6869 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6870 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6872 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6876 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6879bool SPIRVInstructionSelector::selectFirstBitSet32(
6881 unsigned BitSetOpcode)
const {
6882 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6885 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6892bool SPIRVInstructionSelector::selectFirstBitSet64(
6894 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6908 if (ComponentCount > 2) {
6909 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6911 unsigned Opcode) ->
bool {
6912 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6916 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6920 MachineIRBuilder MIRBuilder(
I);
6922 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6926 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6932 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6942 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6943 SPIRV::OpVectorExtractDynamic))
6945 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6946 SPIRV::OpVectorExtractDynamic))
6950 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6951 TII.get(SPIRV::OpVectorShuffle))
6959 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6965 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6966 TII.get(SPIRV::OpVectorShuffle))
6974 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6994 SelectOp = SPIRV::OpSelectSISCond;
6995 AddOp = SPIRV::OpIAddS;
7003 SelectOp = SPIRV::OpSelectVIVCond;
7004 AddOp = SPIRV::OpIAddV;
7010 Register RegSecondaryOffset = Reg0;
7014 if (SwapPrimarySide) {
7015 PrimaryReg = LowReg;
7016 SecondaryReg = HighReg;
7017 RegPrimaryOffset = Reg0;
7018 RegSecondaryOffset = Reg32;
7023 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
7024 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
7029 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
7030 SPIRV::OpINotEqual))
7037 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
7038 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
7043 if (SwapPrimarySide) {
7045 if (!selectOpWithSrcs(RegAdd, ResType,
I,
7046 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
7057 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
7058 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
7063 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
7064 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
7067 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
7071bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
7072 SPIRVTypeInst ResType,
7074 bool IsSigned)
const {
7076 Register OpReg =
I.getOperand(2).getReg();
7079 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
7080 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
7084 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7086 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7088 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7091 return diagnoseUnsupported(
7093 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
7097bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
7098 SPIRVTypeInst ResType,
7099 MachineInstr &
I)
const {
7101 Register OpReg =
I.getOperand(2).getReg();
7106 unsigned ExtendOpcode = SPIRV::OpUConvert;
7107 unsigned BitSetOpcode = GL::FindILsb;
7111 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7113 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7115 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7118 return diagnoseUnsupported(
I,
7119 "spv_firstbitlow only supports 16,32,64 bits.");
7123bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
7124 SPIRVTypeInst ResType,
7125 MachineInstr &
I)
const {
7130 bool UseUntypedPointers =
7131 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7132 unsigned Opcode = UseUntypedPointers
7133 ? SPIRV::OpUntypedVariableLengthArrayINTEL
7134 : SPIRV::OpVariableLengthArrayINTEL;
7136 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
7142 if (UseUntypedPointers) {
7145 "untyped variable length array result must have a recorded element "
7150 MIB.
addUse(
I.getOperand(2).getReg());
7154 unsigned Alignment =
I.getOperand(3).getImm();
7168 while (!Worklist.
empty()) {
7170 switch (
T->getOpcode()) {
7171 case SPIRV::OpTypeInt:
7172 case SPIRV::OpTypeFloat:
7173 case SPIRV::OpTypePointer:
7175 case SPIRV::OpTypeVector:
7176 case SPIRV::OpTypeVectorIdEXT:
7177 case SPIRV::OpTypeMatrix:
7178 case SPIRV::OpTypeArray: {
7179 Register OperandReg =
T->getOperand(1).getReg();
7183 case SPIRV::OpTypeStruct:
7184 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7185 Register OperandReg =
T->getOperand(Idx).getReg();
7200 Register TypeReg = Ty->getOperand(0).getReg();
7201 if (!Visited.
insert(TypeReg).second)
7204 switch (Ty->getOpcode()) {
7205 case SPIRV::OpTypePointer:
7206 if (Ty->getOperand(1).getImm() == SPIRV::StorageClass::StorageBuffer)
7210 case SPIRV::OpTypeArray:
7211 case SPIRV::OpTypeRuntimeArray:
7214 case SPIRV::OpTypeStruct:
7215 for (
unsigned I = 1;
I < Ty->getNumOperands(); ++
I)
7231bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7232 assert(
I.getNumExplicitOperands() == 2);
7234 Register MsgReg =
I.getOperand(1).getReg();
7236 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7239 return diagnoseUnsupported(
7241 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7242 "scalar, pointer, vector, matrix, or aggregate of such types)");
7245 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7252bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7263 uint32_t MsgVal = ~0
u;
7264 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7265 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7268 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7271 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7278bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7279 SPIRVTypeInst ResType,
7280 MachineInstr &
I)
const {
7287 bool UseUntypedPointers =
7288 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7290 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7293 MachineIRBuilder MIRBuilder(
I);
7296 .
addImm(SPIRV::Extension::SPV_KHR_variable_pointers);
7298 .
addImm(SPIRV::Capability::VariablePointersStorageBuffer);
7301 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7304 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7308 if (UseUntypedPointers) {
7312 return diagnoseUnsupported(
7313 I,
"could not deduce the data type of an untyped variable");
7319 unsigned Alignment =
I.getOperand(2).getImm();
7326bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7331 const MachineInstr *PrevI =
I.getPrevNode();
7333 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7337 .
addMBB(
I.getOperand(0).getMBB())
7342 .
addMBB(
I.getOperand(0).getMBB())
7347bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7358 const MachineInstr *NextI =
I.getNextNode();
7360 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7366 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7368 .
addUse(
I.getOperand(0).getReg())
7369 .
addMBB(
I.getOperand(1).getMBB())
7375bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7376 MachineInstr &
I)
const {
7378 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7380 const unsigned NumOps =
I.getNumOperands();
7381 for (
unsigned i = 1; i <
NumOps; i += 2) {
7382 MIB.
addUse(
I.getOperand(i + 0).getReg());
7383 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7389bool SPIRVInstructionSelector::selectGlobalValue(
7390 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7392 MachineIRBuilder MIRBuilder(
I);
7393 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7396 std::string GlobalIdent;
7398 unsigned &
ID = UnnamedGlobalIDs[GV];
7400 ID = UnnamedGlobalIDs.
size();
7401 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7427 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7434 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
7439 MachineInstrBuilder MIB1 =
7440 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7443 MachineInstrBuilder MIB2 =
7445 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7449 GR.
add(ConstVal, MIB2);
7457 MachineInstrBuilder MIB3 =
7458 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7461 GR.
add(ConstVal, MIB3);
7467 assert(NewReg != ResVReg);
7468 return BuildCOPY(ResVReg, NewReg,
I);
7478 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7481 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7487 SPIRVTypeInst ResType =
7491 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7496 if (
GlobalVar->isExternallyInitialized() &&
7497 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7498 constexpr unsigned ReadWriteINTEL = 3u;
7501 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7507bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7508 SPIRVTypeInst ResType,
7509 MachineInstr &
I)
const {
7511 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7519 MachineIRBuilder MIRBuilder(
I);
7524 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7527 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7529 .
add(
I.getOperand(1))
7543 APFloat::rmNearestTiesToEven, &LosesInfo);
7548 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
7558bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7559 SPIRVTypeInst ResType,
7560 MachineInstr &
I)
const {
7563 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7569 Register ExpReg =
I.getOperand(2).getReg();
7571 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7572 SPIRV::OpConvertSToF))
7574 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7581bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7582 SPIRVTypeInst ResType,
7583 MachineInstr &
I)
const {
7599 MachineIRBuilder MIRBuilder(
I);
7600 SPIRVTypeInst FloatType =
7604 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7617 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7618 const bool IsUntyped =
7619 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7621 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7622 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7623 : SPIRV::OpVariable))
7626 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7634 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7637 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7640 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7644 Register IntegralPartReg =
I.getOperand(1).getReg();
7647 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7657 assert(
false &&
"GLSL::Modf is deprecated.");
7668bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7669 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7670 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7671 MachineIRBuilder MIRBuilder(
I);
7672 const SPIRVTypeInst Vec3Ty =
7675 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7687 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7691 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
7697 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7704 assert(
I.getOperand(2).isReg());
7705 const uint32_t ThreadId =
foldImm(
I.getOperand(2),
MRI);
7709 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7720bool SPIRVInstructionSelector::loadBuiltinInputID(
7721 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7722 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7723 MachineIRBuilder MIRBuilder(
I);
7725 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7740 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7744 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7753SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7754 MachineInstr &
I)
const {
7755 MachineIRBuilder MIRBuilder(
I);
7766bool SPIRVInstructionSelector::loadHandleBeforePosition(
7767 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7768 MachineInstr &Pos)
const {
7771 Intrinsic::spv_resource_handlefrombinding);
7779 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7780 MachineIRBuilder MIRBuilder(HandleDef);
7781 SPIRVTypeInst VarType = ResType;
7782 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7784 if (IsStructuredBuffer) {
7793 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7796 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7797 ArraySize, IndexReg, Name, MIRBuilder);
7801 uint32_t LoadOpcode =
7802 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7812bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7813 MachineInstr &
I)
const {
7815 return diagnoseUnsupported(
7816 I,
"this instruction is only supported in shaders.");
7821InstructionSelector *
7825 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
static MaybeAlign getAlign(Value *Ptr)
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)
MachineRegisterInfo & MRI
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 bool containsStorageBufferPointer(SPIRVTypeInst Ty, const SPIRVGlobalRegistry &GR, SmallSet< Register, 8 > &Visited)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef)
static unsigned getPtrCmpOpcode(unsigned Pred)
static void addMemoryOperands(const MachineMemOperand *MemOp, MachineInstrBuilder &MIB, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry &GR, std::optional< Align > AlignOverride=std::nullopt)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
static Register convertPtrToInt(Register Reg, LLT ConvTy, SPIRVTypeInst SpvType, LegalizerHelper &Helper, MachineRegisterInfo &MRI, SPIRVGlobalRegistry *GR)
#define SPIRV_BACKEND_SERVICE_FUN_NAME
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.
LLVM_ABI Align getAlign() const
Return the minimum known alignment in bytes of the actual memory reference.
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 Align[]
Key for Kernel::Arg::Metadata::mAlign.
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.
ElementType
The element type of an SRV or UAV resource.
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)
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