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);
1628 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1635 const bool IsUntyped =
1636 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1638 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1639 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1640 : SPIRV::OpVariable))
1643 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1649 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1652 .
addImm(
static_cast<uint32_t
>(Ex.first))
1654 .
add(
I.getOperand(2))
1658 Register ExpResReg =
I.getOperand(1).getReg();
1660 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1670bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1671 SPIRVTypeInst ResType,
1672 MachineInstr &
I)
const {
1673 Register XReg =
I.getOperand(1).getReg();
1674 Register ExpReg =
I.getOperand(2).getReg();
1680 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1681 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1683 ExpReg = buildVectorSplat(ExpReg, NumElts,
I);
1686 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1687 true,
false, {XReg, ExpReg});
1690bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1691 SPIRVTypeInst ResType,
1692 MachineInstr &
I)
const {
1693 Register CosResVReg =
I.getOperand(1).getReg();
1694 unsigned SrcIdx =
I.getNumExplicitDefs();
1699 MachineIRBuilder MIRBuilder(
I);
1701 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1708 const bool IsUntyped =
1709 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1711 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1712 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1713 : SPIRV::OpVariable))
1716 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1720 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1723 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1725 .
add(
I.getOperand(SrcIdx))
1729 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1737 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1740 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1742 .
add(
I.getOperand(SrcIdx))
1744 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1747 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1749 .
add(
I.getOperand(SrcIdx))
1756bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1757 SPIRVTypeInst ResType,
1760 unsigned Opcode)
const {
1761 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1771bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1772 SPIRVTypeInst ResType,
1775 unsigned Opcode)
const {
1776 MachineIRBuilder MIRBuilder(
I);
1778 Register OpReg =
I.getOperand(1).getReg();
1784 SPIRVTypeInst ExtType =
1791 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1795 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1798 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1801bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1802 SPIRVTypeInst ResType,
1805 unsigned Opcode)
const {
1806 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1809bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1810 SPIRVTypeInst ResType,
1813 unsigned Opcode)
const {
1814 MachineIRBuilder MIRBuilder(
I);
1823 SPIRVTypeInst WorkingType =
1830 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {SrcReg}, SPIRV::OpUConvert))
1834 if (!selectOpWithSrcs(LowCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1842 IsScalar ? SPIRV::OpShiftRightLogicalS : SPIRV::OpShiftRightLogicalV;
1844 if (!selectOpWithSrcs(Shift, SrcType,
I, {SrcReg, ShiftAmount}, ShiftOp))
1848 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {Shift}, SPIRV::OpUConvert))
1852 if (!selectOpWithSrcs(HighCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1857 if (!selectOpWithSrcs(Sum, WorkingType,
I, {HighCount, LowCount},
1858 IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV))
1862 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1863 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1866bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1867 SPIRVTypeInst ResType,
1869 unsigned Opcode)
const {
1874 if (!STI.getTargetTriple().isVulkanOS())
1875 return selectUnOp(ResVReg, ResType,
I, Opcode);
1877 Register OpReg =
I.getOperand(1).getReg();
1880 : SPIRV::OpUConvert;
1884 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1886 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1888 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1890 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1894bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1895 SPIRVTypeInst ResType,
1897 unsigned Opcode)
const {
1899 Register SrcReg =
I.getOperand(1).getReg();
1904 unsigned DefOpCode = DefIt->getOpcode();
1905 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1908 if (
auto *VRD =
getVRegDef(*
MRI, DefIt->getOperand(1).getReg()))
1909 DefOpCode = VRD->getOpcode();
1911 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1912 DefOpCode == TargetOpcode::G_CONSTANT ||
1913 DefOpCode == SPIRV::OpVariable ||
1914 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1915 DefOpCode == SPIRV::OpConstantI) {
1921 uint32_t SpecOpcode = 0;
1923 case SPIRV::OpConvertPtrToU:
1924 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1926 case SPIRV::OpConvertUToPtr:
1927 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1932 TII.get(SPIRV::OpSpecConstantOp))
1942 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1946bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1947 SPIRVTypeInst ResType,
1948 MachineInstr &
I)
const {
1949 Register OpReg =
I.getOperand(1).getReg();
1950 SPIRVTypeInst OpType =
1953 return diagnoseUnsupported(
1954 I,
"incompatible result and operand types in a bitcast");
1955 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1961 std::optional<Align> AlignOverride = std::nullopt) {
1966 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1967 if (
MemOp->isNonTemporal())
1968 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1970 if (!ST->isShader() &&
MemOp->getAlign().value())
1971 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1975 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1976 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1980 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1982 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1986 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1990 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1992 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1993 MIB.
addImm(AlignOverride.value_or(
MemOp->getAlign()).value());
2004 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2006 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2008 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2012bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2013 SPIRVTypeInst ResType,
2014 MachineInstr &
I)
const {
2016 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2021 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2022 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2024 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2026 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2030 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2034 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2035 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2036 I.getDebugLoc(),
I);
2040 MachineIRBuilder MIRBuilder(
I);
2042 if (
I.getNumMemOperands()) {
2043 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2044 if (MemOp->isAtomic())
2045 return selectAtomicLoad(ResVReg, ResType,
I);
2048 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2052 if (!
I.getNumMemOperands()) {
2053 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2055 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2064Register SPIRVInstructionSelector::createPtrSizedIntReg(
2065 MachineIRBuilder &MIRBuilder)
const {
2066 SPIRVTypeInst IntType =
2076SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2077 MachineIRBuilder &MIRBuilder)
const {
2078 SPIRVTypeInst IntType =
2080 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2081 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2089Register SPIRVInstructionSelector::castPtrToPtrToInt(
2090 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2091 MachineIRBuilder &MIRBuilder)
const {
2092 SPIRVTypeInst IntType =
2094 SPIRVTypeInst PtrType =
2108bool SPIRVInstructionSelector::selectAtomicPtrValue(
2109 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2110 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2118 Register IntResult = EmitAtomic(IntType);
2120 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2128bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2129 SPIRVTypeInst ResType,
2130 MachineInstr &
I)
const {
2131 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2134 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2137 return diagnoseUnsupported(
2138 I,
"Lowering to SPIR-V of atomic load is only "
2139 "allowed for integer, floating point or pointer types");
2141 assert(
I.getNumMemOperands());
2142 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2143 assert(MemOp.isAtomic());
2145 uint32_t
Scope =
static_cast<uint32_t
>(
2147 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2153 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2154 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2157 Register MemSemReg = buildI32ConstantInEntryBlock(Sem,
I);
2159 MachineIRBuilder MIRBuilder(
I);
2163 return diagnoseUnsupported(
2164 I,
"Lowering to SPIR-V of atomic load is only "
2165 "allowed for pointer types for physical addressing model");
2170 SPIRV::StorageClass::StorageClass SC =
2172 return selectAtomicPtrValue(
2173 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2174 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2175 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2186 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2197bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2199 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2200 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2205 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2206 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2208 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2213 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2217 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2218 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2219 SPIRVTypeInst SampledType =
2221 SPIRVTypeInst StoreValCompType =
2223 if (StoreValCompType && StoreValCompType != SampledType) {
2226 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2229 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2234 StoreVal = PackedReg;
2237 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2238 TII.get(SPIRV::OpImageWrite))
2244 if (sampledTypeIsSignedInteger(LLVMHandleType))
2247 BMI.constrainAllUses(
TII,
TRI, RBI);
2254 if (PointeeTy && PointeeTy->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
2255 StoreTy->
getOpcode() != SPIRV::OpTypeVectorIdEXT &&
2269 if (
I.getNumMemOperands()) {
2270 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2271 if (MemOp->isAtomic())
2272 return selectAtomicStore(
I);
2279 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2280 PtrSC == SPIRV::StorageClass::Input ||
2281 PtrSC == SPIRV::StorageClass::PushConstant)
2282 return diagnoseUnsupported(
2283 I,
"store into a read-only SPIR-V storage class is not allowed");
2285 MachineIRBuilder MIRBuilder(
I);
2287 if (!
I.getNumMemOperands()) {
2288 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2290 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2299bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2300 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2303 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2304 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2309 if (!PointeeType && PtrType &&
2310 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2313 return diagnoseUnsupported(
I,
2314 "Lowering to SPIR-V of atomic store is only "
2315 "allowed for integer or floating point types");
2317 assert(
I.getNumMemOperands());
2318 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2319 assert(MemOp.isAtomic());
2321 uint32_t
Scope =
static_cast<uint32_t
>(
2323 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2329 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2330 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2333 Register MemSemReg = buildI32ConstantInEntryBlock(Sem,
I);
2334 MachineIRBuilder MIRBuilder(
I);
2338 return diagnoseUnsupported(
2339 I,
"Lowering to SPIR-V of atomic store is only "
2340 "allowed for pointer types for physical addressing model");
2345 SPIRV::StorageClass::StorageClass SC =
2347 return selectAtomicPtrValue(
2348 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2350 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2363 return diagnoseUnsupported(
I,
2364 "Lowering to SPIR-V of atomic store is only "
2365 "allowed for integer or floating point types");
2367 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2377bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2378 SPIRVTypeInst ResType,
2379 MachineInstr &
I)
const {
2380 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2388 const Register PtrsReg =
I.getOperand(2).getReg();
2389 const uint32_t
Alignment =
I.getOperand(3).getImm();
2390 const Register MaskReg =
I.getOperand(4).getReg();
2391 const Register PassthruReg =
I.getOperand(5).getReg();
2392 const Register AlignmentReg = buildI32ConstantInEntryBlock(Alignment,
I);
2396 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2407bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2408 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2415 const Register ValuesReg =
I.getOperand(1).getReg();
2416 const Register PtrsReg =
I.getOperand(2).getReg();
2417 const uint32_t
Alignment =
I.getOperand(3).getImm();
2418 const Register MaskReg =
I.getOperand(4).getReg();
2419 const Register AlignmentReg = buildI32ConstantInEntryBlock(Alignment,
I);
2423 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2432bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2433 const Twine &
Msg)
const {
2434 const Function &
F =
I.getMF()->getFunction();
2435 F.getContext().diagnose(
2436 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2440bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2441 SPIRVTypeInst ResType,
2442 MachineInstr &
I)
const {
2443 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2444 return diagnoseUnsupported(
2445 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2446 "SPIR-V extension: SPV_INTEL_variable_length_array");
2448 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2455bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2456 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2457 return diagnoseUnsupported(
2459 "llvm.stackrestore intrinsic: this instruction requires the following "
2460 "SPIR-V extension: SPV_INTEL_variable_length_array");
2461 if (!
I.getOperand(0).isReg())
2464 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2465 .
addUse(
I.getOperand(0).getReg())
2471SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2472 MachineIRBuilder MIRBuilder(
I);
2473 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2480 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2484 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2485 Type *ArrTy = ArrayType::get(ValTy, Num);
2487 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2490 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2501 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2502 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2503 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2504 : SPIRV::OpVariable))
2507 .
addImm(SPIRV::StorageClass::UniformConstant);
2520bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2523 Register DstReg =
I.getOperand(0).getReg();
2529 return diagnoseUnsupported(
2530 I,
"OpCopyMemory requires operands to have the same type");
2535 return diagnoseUnsupported(
2536 I,
"Unable to determine pointee type size for OpCopyMemory");
2537 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2538 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2539 return diagnoseUnsupported(
2540 I,
"OpCopyMemory requires the size to match the pointee type size");
2543 const unsigned Opcode =
2544 IsLogical ? SPIRV::OpCopyMemory : SPIRV::OpCopyMemorySized;
2545 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2549 MIB.
addUse(
I.getOperand(2).getReg());
2551 if (
I.getNumMemOperands()) {
2552 MachineIRBuilder MIRBuilder(
I);
2553 const MachineMemOperand *DstMemOp = *
I.memoperands_begin();
2555 Align SrcAlign = DstAlign;
2558 if (
I.getNumMemOperands() > 1 && !STI.
isShader())
2559 SrcAlign = (*std::next(
I.memoperands_begin()))->
getAlign();
2569 std::min(DstAlign, SrcAlign));
2576bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2577 MachineInstr &
I)
const {
2579 Register SizeReg =
I.getOperand(2).getReg();
2581 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2585 Register SrcReg =
I.getOperand(1).getReg();
2586 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2587 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2588 Register VarReg = getOrCreateMemSetGlobal(
I);
2591 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2593 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2595 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2598 if (!selectCopyMemory(
I, SrcReg))
2600 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2601 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2606bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2607 SPIRVTypeInst ResType,
2610 unsigned NegateOpcode)
const {
2612 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2613 uint32_t
Scope =
static_cast<uint32_t
>(
2615 MemOp->getSyncScopeID()));
2616 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2618 Register Ptr =
I.getOperand(1).getReg();
2619 uint32_t ScSem =
static_cast<uint32_t
>(
2623 Register MemSemReg = buildI32ConstantInEntryBlock(
2627 Register ValueReg =
I.getOperand(2).getReg();
2628 if (NegateOpcode != 0) {
2631 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2637 if (NewOpcode != SPIRV::OpAtomicExchange)
2638 return diagnoseUnsupported(
2639 I,
"Lowering to SPIR-V of this atomic operation is not "
2640 "allowed for pointer types");
2642 return diagnoseUnsupported(
2643 I,
"Lowering to SPIR-V of atomic exchange is only "
2644 "allowed for pointer types for physical addressing model");
2651 MachineIRBuilder MIRBuilder(
I);
2653 return selectAtomicPtrValue(
2654 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2656 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2657 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2658 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2666 return ExchangeResReg;
2670 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2681bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2682 unsigned ArgI =
I.getNumOperands() - 1;
2684 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2685 SPIRVTypeInst SrcType =
2689 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2693 unsigned CurrentIndex = 0;
2694 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2695 Register ResVReg =
I.getOperand(i).getReg();
2704 ResType = ScalarType;
2713 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2719 for (
unsigned j = 0;
j < NumElements; ++
j) {
2720 MIB.
addImm(CurrentIndex + j);
2722 CurrentIndex += NumElements;
2726 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2738bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2741 ? SPIRV::MemorySemantics::UniformMemory |
2742 SPIRV::MemorySemantics::WorkgroupMemory |
2743 SPIRV::MemorySemantics::ImageMemory
2744 : SPIRV::MemorySemantics::WorkgroupMemory |
2745 SPIRV::MemorySemantics::CrossWorkgroupMemory |
2746 SPIRV::MemorySemantics::ImageMemory;
2748 STI.getTargetTriple(),
static_cast<uint32_t
>(
getMemSemantics(AO)), ScSem);
2749 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
2753 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2755 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2762bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2763 SPIRVTypeInst ResType,
2765 unsigned Opcode)
const {
2766 Type *ResTy =
nullptr;
2769 return diagnoseUnsupported(
2771 "Not enough info to select the arithmetic with overflow instruction");
2773 return diagnoseUnsupported(
I,
2774 "Expect struct type result for the arithmetic "
2775 "with overflow instruction");
2781 MachineIRBuilder MIRBuilder(
I);
2783 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2784 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2791 Register ZeroReg = buildZerosVal(ResType,
I);
2796 if (ResName.
size() > 0)
2804 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2805 MIB.
addUse(
I.getOperand(i).getReg());
2811 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2813 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2814 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2821 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2822 .
addDef(
I.getOperand(1).getReg())
2830bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2831 SPIRVTypeInst ResType,
2832 MachineInstr &
I)
const {
2834 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2835 Register Ptr =
I.getOperand(2).getReg();
2836 Register ScopeReg =
I.getOperand(5).getReg();
2837 Register MemSemEqReg =
I.getOperand(6).getReg();
2838 Register MemSemNeqReg =
I.getOperand(7).getReg();
2840 Register Val =
I.getOperand(4).getReg();
2844 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2863 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2870 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2882 case SPIRV::StorageClass::DeviceOnlyINTEL:
2883 case SPIRV::StorageClass::HostOnlyINTEL:
2892 bool IsGRef =
false;
2893 bool IsAllowedRefs =
2894 llvm::all_of(
MRI->use_instructions(ResVReg), [&IsGRef](
auto const &It) {
2895 unsigned Opcode = It.getOpcode();
2896 if (Opcode == SPIRV::OpConstantComposite ||
2897 Opcode == SPIRV::OpSpecConstantComposite ||
2898 Opcode == SPIRV::OpVariable ||
2899 Opcode == SPIRV::OpUntypedVariableKHR ||
2900 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2901 return IsGRef = true;
2902 return Opcode == SPIRV::OpName;
2904 return IsAllowedRefs && IsGRef;
2907Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2908 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2910 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2914SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2916 uint32_t Opcode)
const {
2917 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2918 TII.get(SPIRV::OpSpecConstantOp))
2926SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2927 SPIRVTypeInst SrcPtrTy)
const {
2928 SPIRVTypeInst GenericPtrTy =
2932 SPIRV::StorageClass::Generic),
2936 MachineInstrBuilder MIB = buildSpecConstantOp(
2938 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2948bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2949 SPIRVTypeInst ResType,
2950 MachineInstr &
I)
const {
2954 Register SrcPtr =
I.getOperand(1).getReg();
2959 return BuildCOPY(ResVReg, SrcPtr,
I);
2969 unsigned SpecOpcode = [&]() ->
unsigned {
2970 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2971 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2973 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2975 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
2983 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2985 .constrainAllUses(
TII,
TRI, RBI);
2987 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2989 buildSpecConstantOp(
2991 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2992 .constrainAllUses(
TII,
TRI, RBI);
2999 return BuildCOPY(ResVReg, SrcPtr,
I);
3001 if ((SrcSC == SPIRV::StorageClass::Function &&
3002 DstSC == SPIRV::StorageClass::Private) ||
3003 (DstSC == SPIRV::StorageClass::Function &&
3004 SrcSC == SPIRV::StorageClass::Private))
3005 return BuildCOPY(ResVReg, SrcPtr,
I);
3009 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3012 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3015 SPIRVTypeInst GenericPtrTy =
3034 return selectUnOp(ResVReg, ResType,
I,
3035 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
3037 return selectUnOp(ResVReg, ResType,
I,
3038 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
3040 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3042 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3052bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3053 SPIRVTypeInst ResType,
3054 MachineInstr &
I)
const {
3056 return diagnoseUnsupported(
3057 I,
"G_PTRMASK is not supported with logical SPIR-V");
3062 Register PtrReg =
I.getOperand(1).getReg();
3063 Register MaskReg =
I.getOperand(2).getReg();
3082 ? SPIRV::OpBitwiseAndV
3083 : SPIRV::OpBitwiseAndS;
3106 return SPIRV::OpFOrdEqual;
3108 return SPIRV::OpFOrdGreaterThanEqual;
3110 return SPIRV::OpFOrdGreaterThan;
3112 return SPIRV::OpFOrdLessThanEqual;
3114 return SPIRV::OpFOrdLessThan;
3116 return SPIRV::OpFOrdNotEqual;
3118 return SPIRV::OpOrdered;
3120 return SPIRV::OpFUnordEqual;
3122 return SPIRV::OpFUnordGreaterThanEqual;
3124 return SPIRV::OpFUnordGreaterThan;
3126 return SPIRV::OpFUnordLessThanEqual;
3128 return SPIRV::OpFUnordLessThan;
3130 return SPIRV::OpFUnordNotEqual;
3132 return SPIRV::OpUnordered;
3142 return SPIRV::OpIEqual;
3144 return SPIRV::OpINotEqual;
3146 return SPIRV::OpSGreaterThanEqual;
3148 return SPIRV::OpSGreaterThan;
3150 return SPIRV::OpSLessThanEqual;
3152 return SPIRV::OpSLessThan;
3154 return SPIRV::OpUGreaterThanEqual;
3156 return SPIRV::OpUGreaterThan;
3158 return SPIRV::OpULessThanEqual;
3160 return SPIRV::OpULessThan;
3169 return SPIRV::OpPtrEqual;
3171 return SPIRV::OpPtrNotEqual;
3182 return SPIRV::OpLogicalEqual;
3184 return SPIRV::OpLogicalNotEqual;
3222bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3223 SPIRVTypeInst ResType,
3225 unsigned OpAnyOrAll)
const {
3226 assert(
I.getNumOperands() == 3);
3227 assert(
I.getOperand(2).isReg());
3229 Register InputRegister =
I.getOperand(2).getReg();
3232 assert(InputType &&
"VReg has no type assigned");
3236 assert(ResVReg ==
I.getOperand(0).getReg());
3237 return BuildCOPY(ResVReg, InputRegister,
I);
3241 unsigned SpirvNotEqualId =
3242 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3244 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3249 IsBoolTy ? InputRegister
3257 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3259 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3276bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3277 SPIRVTypeInst ResType,
3278 MachineInstr &
I)
const {
3279 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3282bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3283 SPIRVTypeInst ResType,
3284 MachineInstr &
I)
const {
3285 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3289bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3290 SPIRVTypeInst ResType,
3291 MachineInstr &
I)
const {
3292 assert(
I.getNumOperands() == 4);
3293 assert(
I.getOperand(2).isReg());
3294 assert(
I.getOperand(3).isReg());
3296 [[maybe_unused]] SPIRVTypeInst VecType =
3301 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3302 "dot product requires either a vector of at least 2 components or"
3303 " the SPV_EXT_long vector extension.");
3305 [[maybe_unused]] SPIRVTypeInst EltType =
3314 .
addUse(
I.getOperand(2).getReg())
3315 .
addUse(
I.getOperand(3).getReg())
3320bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3321 SPIRVTypeInst ResType,
3324 assert(
I.getNumOperands() == 4);
3325 assert(
I.getOperand(2).isReg());
3326 assert(
I.getOperand(3).isReg());
3329 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3333 .
addUse(
I.getOperand(2).getReg())
3334 .
addUse(
I.getOperand(3).getReg())
3341bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3342 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3343 assert(
I.getNumOperands() == 4);
3344 assert(
I.getOperand(2).isReg());
3345 assert(
I.getOperand(3).isReg());
3349 Register Vec0 =
I.getOperand(2).getReg();
3350 Register Vec1 =
I.getOperand(3).getReg();
3354 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3363 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3364 "dot product requires either a vector of at least 2 components "
3365 "or the SPV_EXT_long_vector extension.");
3368 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3378 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3389 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3401bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3402 SPIRVTypeInst ResType,
3403 MachineInstr &
I)
const {
3405 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3408 .
addUse(
I.getOperand(2).getReg())
3413bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3414 SPIRVTypeInst ResType,
3415 MachineInstr &
I)
const {
3417 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3420 .
addUse(
I.getOperand(2).getReg())
3425bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3426 SPIRVTypeInst ResType,
3427 MachineInstr &
I)
const {
3430 Register Src =
I.getOperand(2).getReg();
3473bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3474 SPIRVTypeInst ResType,
3475 MachineInstr &
I)
const {
3477 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3480 .
addUse(
I.getOperand(2).getReg())
3485template <
bool Signed>
3486bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3487 SPIRVTypeInst ResType,
3488 MachineInstr &
I)
const {
3489 assert(
I.getNumOperands() == 5);
3490 assert(
I.getOperand(2).isReg());
3491 assert(
I.getOperand(3).isReg());
3492 assert(
I.getOperand(4).isReg());
3495 Register Acc =
I.getOperand(2).getReg();
3499 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3501 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3506 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3509 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3521template <
bool Signed>
3522bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3523 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3524 assert(
I.getNumOperands() == 5);
3525 assert(
I.getOperand(2).isReg());
3526 assert(
I.getOperand(3).isReg());
3527 assert(
I.getOperand(4).isReg());
3530 Register Acc =
I.getOperand(2).getReg();
3536 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3540 for (
unsigned i = 0; i < 4; i++) {
3563 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3583 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3598bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3599 SPIRVTypeInst ResType,
3600 MachineInstr &
I)
const {
3601 assert(
I.getNumOperands() == 3);
3602 assert(
I.getOperand(2).isReg());
3604 Register VZero = buildZerosValF(ResType,
I);
3605 Register VOne = buildOnesValF(ResType,
I);
3607 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3610 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3612 .
addUse(
I.getOperand(2).getReg())
3619bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3620 SPIRVTypeInst ResType,
3621 MachineInstr &
I)
const {
3622 assert(
I.getNumOperands() == 3);
3623 assert(
I.getOperand(2).isReg());
3625 Register InputRegister =
I.getOperand(2).getReg();
3627 auto &
DL =
I.getDebugLoc();
3630 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3637 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3639 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3647 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3652 if (NeedsConversion) {
3653 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3664bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3665 SPIRVTypeInst ResType,
3667 unsigned Opcode)
const {
3671 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3677 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3678 BMI.addUse(
I.getOperand(J).getReg());
3685bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3688 bool WithGroupSync)
const {
3690 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3692 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3694 assert(((Scope != SPIRV::Scope::Workgroup) ||
3695 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3696 "Workgroup Scope must set WorkGroupMemory semantic "
3697 "in Barrier instruction");
3699 assert(((Scope != SPIRV::Scope::Device) ||
3700 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3701 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3702 "Device Scope must set UniformMemory and ImageMemory semantic "
3703 "in Barrier instruction");
3709 if (WithGroupSync) {
3710 Register ExecReg = buildI32ConstantInEntryBlock(SPIRV::Scope::Workgroup,
I);
3714 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
3715 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
3717 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3721bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3722 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3727 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3728 SPIRV::OpGroupNonUniformBallot))
3733 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3738 .
addImm(SPIRV::GroupOperation::Reduce)
3745bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3746 SPIRVTypeInst ResType,
3747 MachineInstr &
I)
const {
3752 Register InputReg =
I.getOperand(2).getReg();
3757 bool IsVector = NumElems > 1 ||
3758 (InputType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
3772 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3773 SPIRV::OpGroupNonUniformAllEqual);
3778 ElementResults.
reserve(NumElems);
3780 for (
unsigned Idx = 0;
Idx < NumElems; ++
Idx) {
3793 ElemInput = Extracted;
3799 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3810 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3821bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3822 SPIRVTypeInst ResType,
3823 MachineInstr &
I)
const {
3825 assert(
I.getNumOperands() == 3);
3827 auto Op =
I.getOperand(2);
3837 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3839 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3840 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3861 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3865 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3872bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3873 SPIRVTypeInst ResType,
3875 bool IsUnsigned)
const {
3876 return selectWaveReduce(
3877 ResVReg, ResType,
I, IsUnsigned,
3878 [&](
Register InputRegister,
bool IsUnsigned) {
3879 const bool IsFloatTy =
3881 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3882 : SPIRV::OpGroupNonUniformSMax;
3883 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3887bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3888 SPIRVTypeInst ResType,
3890 bool IsUnsigned)
const {
3891 return selectWaveReduce(
3892 ResVReg, ResType,
I, IsUnsigned,
3893 [&](
Register InputRegister,
bool IsUnsigned) {
3894 const bool IsFloatTy =
3896 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3897 : SPIRV::OpGroupNonUniformSMin;
3898 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3902bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3903 SPIRVTypeInst ResType,
3904 MachineInstr &
I)
const {
3905 return selectWaveReduce(ResVReg, ResType,
I,
false,
3906 [&](
Register InputRegister,
bool IsUnsigned) {
3908 InputRegister, SPIRV::OpTypeFloat);
3909 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3910 : SPIRV::OpGroupNonUniformIAdd;
3914bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3915 SPIRVTypeInst ResType,
3916 MachineInstr &
I)
const {
3917 return selectWaveReduce(ResVReg, ResType,
I,
false,
3918 [&](
Register InputRegister,
bool IsUnsigned) {
3920 InputRegister, SPIRV::OpTypeFloat);
3921 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3922 : SPIRV::OpGroupNonUniformIMul;
3926template <
typename PickOpcodeFn>
3927bool SPIRVInstructionSelector::selectWaveReduce(
3928 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3929 PickOpcodeFn &&PickOpcode)
const {
3930 assert(
I.getNumOperands() == 3);
3931 assert(
I.getOperand(2).isReg());
3933 Register InputRegister =
I.getOperand(2).getReg();
3937 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3940 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3946 .
addImm(SPIRV::GroupOperation::Reduce)
3947 .
addUse(
I.getOperand(2).getReg())
3952bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3953 SPIRVTypeInst ResType,
3955 unsigned Opcode)
const {
3956 return selectWaveReduce(
3957 ResVReg, ResType,
I,
false,
3958 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3961bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3962 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3963 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3964 [&](
Register InputRegister,
bool IsUnsigned) {
3966 InputRegister, SPIRV::OpTypeFloat);
3968 ? SPIRV::OpGroupNonUniformFAdd
3969 : SPIRV::OpGroupNonUniformIAdd;
3973bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3974 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3975 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3976 [&](
Register InputRegister,
bool IsUnsigned) {
3978 InputRegister, SPIRV::OpTypeFloat);
3980 ? SPIRV::OpGroupNonUniformFMul
3981 : SPIRV::OpGroupNonUniformIMul;
3985template <
typename PickOpcodeFn>
3986bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3987 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3988 PickOpcodeFn &&PickOpcode)
const {
3989 assert(
I.getNumOperands() == 3);
3990 assert(
I.getOperand(2).isReg());
3992 Register InputRegister =
I.getOperand(2).getReg();
3996 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3999 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
4005 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
4006 .
addUse(
I.getOperand(2).getReg())
4011bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
4012 SPIRVTypeInst ResType,
4015 assert(
I.getNumOperands() == 3);
4016 assert(
I.getOperand(2).isReg());
4018 Register InputRegister =
I.getOperand(2).getReg();
4024 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
4035bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
4036 SPIRVTypeInst ResType,
4043 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
4048 : SPIRV::OpUConvert;
4050 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4053 ShiftOp = SPIRV::OpShiftRightLogicalV;
4058 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4059 TII.get(SPIRV::OpConstantComposite))
4062 for (
unsigned It = 0; It <
N; ++It)
4066 ShiftConst = CompositeReg;
4071 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
4076 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
4081 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
4086 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
4089bool SPIRVInstructionSelector::handle64BitOverflow(
4091 unsigned int Opcode,
4098 "handle64BitOverflow should only be used for integer types");
4100 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4102 MachineIRBuilder MIRBuilder(
I);
4104 SPIRVTypeInst I64x2Type =
4106 SPIRVTypeInst Vec2ResType =
4109 std::vector<Register> PartialRegs;
4111 unsigned CurrentComponent = 0;
4112 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4116 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4117 TII.get(SPIRV::OpVectorShuffle))
4122 .
addImm(CurrentComponent)
4123 .
addImm(CurrentComponent + 1);
4133 PartialRegs.push_back(SubVecReg);
4136 if (CurrentComponent != ComponentCount) {
4142 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4143 SPIRV::OpVectorExtractDynamic))
4152 PartialRegs.push_back(FinalElemResReg);
4156 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4157 SPIRV::OpCompositeConstruct);
4160bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4161 SPIRVTypeInst ResType,
4165 if (ComponentCount > 2)
4166 return handle64BitOverflow(
4167 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4169 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4171 MachineIRBuilder MIRBuilder(
I);
4175 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4179 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4184 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4191 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4192 TII.get(SPIRV::OpVectorShuffle))
4197 for (
unsigned J = 0; J < ComponentCount; ++J) {
4204 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4207bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4208 SPIRVTypeInst ResType,
4212 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4220bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4221 SPIRVTypeInst ResType,
4222 MachineInstr &
I)
const {
4223 Register OpReg =
I.getOperand(1).getReg();
4232 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4234 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4236 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4238 return SPIRVInstructionSelector::diagnoseUnsupported(
4239 I,
"G_BITREVERSE only support 16,32,64 bits.");
4243 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4254 unsigned AndOp = SPIRV::OpBitwiseAndS;
4255 unsigned OrOp = SPIRV::OpBitwiseOrS;
4256 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4257 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4258 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4260 AndOp = SPIRV::OpBitwiseAndV;
4261 OrOp = SPIRV::OpBitwiseOrV;
4262 ShlOp = SPIRV::OpShiftLeftLogicalV;
4263 ShrOp = SPIRV::OpShiftRightLogicalV;
4269 const unsigned Shift) ->
Register {
4272 (ResType->
getOpcode() != SPIRV::OpTypeVectorIdEXT ||
4279 Register MaskReg = CreateConst(Mask);
4280 Register ShiftReg = CreateConst(Shift);
4287 if (!selectOpWithSrcs(
T1, ResType,
I, {
Input, ShiftReg}, ShrOp) ||
4288 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4289 !selectOpWithSrcs(T3, ResType,
I, {
Input, MaskReg}, AndOp) ||
4290 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4291 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4300 while ((Shift >>= 1) > 0) {
4307 return BuildCOPY(ResVReg, Result,
I);
4310bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4311 SPIRVTypeInst ResType,
4312 MachineInstr &
I)
const {
4313 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4314 "G_FREEZE must define and use a register");
4315 Register OpReg =
I.getOperand(1).getReg();
4319 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4333 if (
Def->getOpcode() == TargetOpcode::COPY)
4336 switch (
Def->getOpcode()) {
4337 case SPIRV::ASSIGN_TYPE:
4338 if (MachineInstr *AssignToDef =
4340 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4341 Reg =
Def->getOperand(2).getReg();
4344 case SPIRV::OpUndef:
4345 Reg =
Def->getOperand(1).getReg();
4348 unsigned DestOpCode;
4350 DestOpCode = SPIRV::OpConstantNull;
4351 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4352 "static undef/poison lowered to OpConstantNull\n");
4354 DestOpCode = TargetOpcode::COPY;
4356 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4357 "skipped, lowered as a copy of the operand\n");
4359 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4360 .
addDef(
I.getOperand(0).getReg())
4368bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4369 SPIRVTypeInst ResType,
4370 MachineInstr &
I)
const {
4374 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4378 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4383 for (
unsigned i =
I.getNumExplicitDefs();
4384 i <
I.getNumExplicitOperands() && IsConst; ++i)
4389 return diagnoseUnsupported(
4390 I,
"There must be at least two constituent operands in a vector");
4395 for (
unsigned i =
I.getNumExplicitDefs();
4396 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4402 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4409 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4410 TII.get(IsConst ? SPIRV::OpConstantComposite
4411 : SPIRV::OpCompositeConstruct))
4414 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4415 MIB.
addUse(
I.getOperand(i).getReg());
4420bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4421 SPIRVTypeInst ResType,
4422 MachineInstr &
I)
const {
4426 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4431 unsigned OpIdx =
I.getNumExplicitDefs();
4432 if (!
I.getOperand(OpIdx).isReg())
4436 Register OpReg =
I.getOperand(OpIdx).getReg();
4440 return diagnoseUnsupported(
4441 I,
"There must be at least two constituent operands in a vector");
4444 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4445 TII.get(IsConst ? SPIRV::OpConstantComposite
4446 : SPIRV::OpCompositeConstruct))
4449 for (
unsigned i = 0; i <
N; ++i)
4455bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4456 SPIRVTypeInst ResType,
4457 MachineInstr &
I)
const {
4463 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4465 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4466 TII.get(SPIRV::OpCompositeConstruct))
4469 for (
unsigned OpIdx =
I.getNumExplicitDefs();
4470 OpIdx <
I.getNumExplicitOperands(); ++OpIdx)
4471 MIB.
addUse(
I.getOperand(OpIdx).getReg());
4476bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4477 SPIRVTypeInst ResType,
4478 MachineInstr &
I)
const {
4484 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4486 Opcode = SPIRV::OpDemoteToHelperInvocation;
4488 Opcode = SPIRV::OpKill;
4493 ToErase.eraseFromParent();
4502bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4503 SPIRVTypeInst ResType,
unsigned CmpOpc,
4504 MachineInstr &
I)
const {
4505 Register Cmp0 =
I.getOperand(2).getReg();
4506 Register Cmp1 =
I.getOperand(3).getReg();
4509 "CMP operands should have the same type");
4510 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4520bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4521 SPIRVTypeInst ResType,
4522 MachineInstr &
I)
const {
4523 auto Pred =
I.getOperand(1).getPredicate();
4526 Register CmpOperand =
I.getOperand(2).getReg();
4528 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4533 Register Op1 =
I.getOperand(3).getReg();
4537 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4542 I.getOperand(3).setReg(NewOp1);
4548 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4553Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4554 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4556 SPIRVTypeInst SpvI32Ty =
4558 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4563 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4564 MachineInstr *
MI =
nullptr;
4568 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4572 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4573 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4579 GR.
add(ConstInt,
MI);
4584bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4585 SPIRVTypeInst ResType,
4586 MachineInstr &
I)
const {
4587 auto Pred =
I.getOperand(1).getPredicate();
4595 for (
unsigned Idx = 0;
Idx < 2; ++
Idx) {
4597 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
4600 .
addUse(
I.getOperand(2 + Idx).getReg())
4607 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLogicalOr))
4614 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLogicalNot))
4622 return selectCmp(ResVReg, ResType, CmpOp,
I);
4625bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4626 SPIRVTypeInst ResType,
4627 MachineInstr &
I)
const {
4629 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4639 MachineIRBuilder MIRBuilder(
I);
4646 APFloat ConstVal(3.3219280948873623);
4650 APFloat::rmNearestTiesToEven, &LosesInfo);
4655 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
4657 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4658 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4660 if (!selectExtInst(ResVReg, ResType,
I,
4661 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4671Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4672 MachineInstr &
I)
const {
4680bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4686 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4694 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4697 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4698 Def->getOpcode() == SPIRV::OpConstantI)
4711 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4712 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4714 Intrinsic::spv_const_composite)) {
4715 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4716 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4717 if (!IsZero(
Def->getOperand(i).getReg()))
4726Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4727 MachineInstr &
I)
const {
4736Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4737 MachineInstr &
I)
const {
4747 SPIRVTypeInst ResType,
4748 MachineInstr &
I)
const {
4759 MachineInstr &
I)
const {
4764 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4765 TII.get(SPIRV::OpCompositeConstruct))
4768 for (
unsigned J = 0; J < NumElts; ++J)
4774bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4775 SPIRVTypeInst ResType,
4776 MachineInstr &
I)
const {
4777 Register SelectFirstArg =
I.getOperand(2).getReg();
4778 Register SelectSecondArg =
I.getOperand(3).getReg();
4787 Register CondReg =
I.getOperand(1).getReg();
4788 bool IsScalarBool = GR.
isScalarOfType(CondReg, SPIRV::OpTypeBool);
4796 CondReg = buildVectorSplat(CondReg, NumElts,
I);
4797 IsScalarBool =
false;
4800 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4801 }
else if (IsPtrTy) {
4802 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4804 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4807 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4808 "boolean condition");
4810 Opcode = SPIRV::OpSelectSFSCond;
4811 }
else if (IsPtrTy) {
4812 Opcode = SPIRV::OpSelectSPSCond;
4814 Opcode = SPIRV::OpSelectSISCond;
4817 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4829bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4830 SPIRVTypeInst ResType,
4832 MachineInstr &InsertAt,
4833 bool IsSigned)
const {
4835 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4836 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4837 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4839 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4851bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4852 SPIRVTypeInst ResType,
4853 MachineInstr &
I,
bool IsSigned,
4854 unsigned Opcode)
const {
4855 Register SrcReg =
I.getOperand(1).getReg();
4866 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4868 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4871bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4872 SPIRVTypeInst ResType, MachineInstr &
I,
4873 bool IsSigned)
const {
4874 Register SrcReg =
I.getOperand(1).getReg();
4876 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4880 if (ResType == SrcType)
4881 return BuildCOPY(ResVReg, SrcReg,
I);
4883 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4884 return selectUnOp(ResVReg, ResType,
I, Opcode);
4887bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4888 SPIRVTypeInst ResType,
4890 bool IsSigned)
const {
4891 MachineIRBuilder MIRBuilder(
I);
4892 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
4897 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4905 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4908 .
addUse(
I.getOperand(1).getReg())
4909 .
addUse(
I.getOperand(2).getReg())
4914 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4917 .
addUse(
I.getOperand(1).getReg())
4918 .
addUse(
I.getOperand(2).getReg())
4926 unsigned SelectOpcode =
4927 (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4929 ? SPIRV::OpSelectVIVCond
4930 : SPIRV::OpSelectSISCond;
4935 .
addUse(buildOnesVal(
true, ResType,
I))
4936 .
addUse(buildZerosVal(ResType,
I))
4943 .
addUse(buildOnesVal(
false, ResType,
I))
4948bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4951 SPIRVTypeInst IntTy,
4952 SPIRVTypeInst BoolTy)
const {
4956 isVectorType(IntTy) ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4958 Register One = buildOnesVal(
false, IntTy,
I);
4966 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4975bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4976 SPIRVTypeInst ResType,
4977 MachineInstr &
I)
const {
4978 Register IntReg =
I.getOperand(1).getReg();
4981 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4982 if (ArgType == ResType)
4983 return BuildCOPY(ResVReg, IntReg,
I);
4985 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4986 return selectUnOp(ResVReg, ResType,
I, Opcode);
4989bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4990 SPIRVTypeInst ResType,
4991 MachineInstr &
I)
const {
4992 unsigned Opcode =
I.getOpcode();
4993 unsigned TpOpcode = ResType->
getOpcode();
4995 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
4996 assert(Opcode == TargetOpcode::G_CONSTANT &&
4997 I.getOperand(1).getCImm()->isZero());
4998 MachineBasicBlock &DepMBB =
I.getMF()->front();
5001 }
else if (TpOpcode == SPIRV::OpTypeVectorIdEXT) {
5006 "Expected <1 x T> Vector!");
5007 if (Opcode == TargetOpcode::G_FCONSTANT)
5013 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
5021 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
5024bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
5025 SPIRVTypeInst ResType,
5026 MachineInstr &
I)
const {
5027 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5034bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
5035 SPIRVTypeInst ResType,
5036 MachineInstr &
I)
const {
5038 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
5042 .
addUse(
I.getOperand(3).getReg())
5044 .
addUse(
I.getOperand(2).getReg());
5045 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
5051bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
5052 SPIRVTypeInst ResType,
5053 MachineInstr &
I)
const {
5054 Type *MaybeResTy =
nullptr;
5059 "Expected aggregate type for extractv instruction");
5061 SPIRV::AccessQualifier::ReadWrite,
false);
5065 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5068 .
addUse(
I.getOperand(2).getReg());
5069 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
5075bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
5076 SPIRVTypeInst ResType,
5077 MachineInstr &
I)
const {
5079 return selectInsertVal(ResVReg, ResType,
I);
5081 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
5084 .
addUse(
I.getOperand(2).getReg())
5085 .
addUse(
I.getOperand(3).getReg())
5086 .
addUse(
I.getOperand(4).getReg())
5091bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
5092 SPIRVTypeInst ResType,
5093 MachineInstr &
I)
const {
5095 return selectExtractVal(ResVReg, ResType,
I);
5097 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
5100 .
addUse(
I.getOperand(2).getReg())
5101 .
addUse(
I.getOperand(3).getReg())
5106bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
5107 SPIRVTypeInst ResType,
5108 MachineInstr &
I)
const {
5109 const bool IsGEPInBounds =
I.getOperand(2).getImm();
5112 const bool UseUntypedPointers =
5113 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
5118 if (UseUntypedPointers) {
5120 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
5121 : SPIRV::OpUntypedAccessChainKHR;
5123 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
5124 : SPIRV::OpUntypedPtrAccessChainKHR;
5133 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
5135 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
5136 : SPIRV::OpPtrAccessChain;
5141 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5146 if (UseUntypedPointers) {
5161 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5162 Def->getOperand(1).isReg())
5164 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5165 if (
const auto *GVar =
5168 SPIRV::AccessQualifier::ReadWrite,
5172 return diagnoseUnsupported(
5173 I,
"could not deduce the base type of an untyped access chain");
5178 Res.addUse(BaseReg);
5180 const bool IsAccessChainOpcode =
5181 (Opcode == SPIRV::OpAccessChain ||
5182 Opcode == SPIRV::OpInBoundsAccessChain ||
5183 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5184 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5188 "Cannot translate GEP to OpAccessChain.");
5191 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5192 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5193 Res.addUse(
I.getOperand(i).getReg());
5194 Res.constrainAllUses(
TII,
TRI, RBI);
5203 if (Extract.
getOpcode() == SPIRV::OpCompositeExtract) {
5209 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5215 TII.get(SPIRV::OpCompositeInsert))
5229bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5231 unsigned Lim =
I.getNumExplicitOperands();
5232 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5233 Register OpReg =
I.getOperand(i).getReg();
5237 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5238 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5251 SPIRVTypeInst WrapType = OpType;
5252 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5254 SPIRV::StorageClass::CodeSectionINTEL) {
5256 SPIRV::StorageClass::Function,
I);
5263 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5264 TII.get(SPIRV::OpSpecConstantOp))
5267 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5269 GR.
add(OpDefine, MIB);
5275bool SPIRVInstructionSelector::selectDerivativeInst(
5276 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5277 const unsigned DPdOpCode)
const {
5280 if (!errorIfInstrOutsideShader(
I))
5286 Register SrcReg =
I.getOperand(2).getReg();
5291 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5294 .
addUse(
I.getOperand(2).getReg());
5296 MachineIRBuilder MIRBuilder(
I);
5299 if (componentCount != 1)
5307 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5312 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5317 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5325bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5326 SPIRVTypeInst ResType,
5327 MachineInstr &
I)
const {
5331 case Intrinsic::spv_load:
5332 return selectLoad(ResVReg, ResType,
I);
5333 case Intrinsic::spv_atomic_load:
5334 return selectAtomicLoad(ResVReg, ResType,
I);
5335 case Intrinsic::spv_store:
5336 return selectStore(
I);
5337 case Intrinsic::spv_atomic_store:
5338 return selectAtomicStore(
I);
5339 case Intrinsic::spv_extractv:
5340 return selectExtractVal(ResVReg, ResType,
I);
5341 case Intrinsic::spv_insertv:
5342 return selectInsertVal(ResVReg, ResType,
I);
5343 case Intrinsic::spv_extractelt:
5344 return selectExtractElt(ResVReg, ResType,
I);
5345 case Intrinsic::spv_insertelt:
5346 return selectInsertElt(ResVReg, ResType,
I);
5347 case Intrinsic::spv_gep:
5348 return selectGEP(ResVReg, ResType,
I);
5349 case Intrinsic::spv_bitcast: {
5350 Register OpReg =
I.getOperand(2).getReg();
5351 SPIRVTypeInst OpType =
5355 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5357 case Intrinsic::spv_unref_global:
5358 case Intrinsic::spv_init_global: {
5364 Register GVarVReg =
MI->getOperand(0).getReg();
5365 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5370 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5372 MI->eraseFromParent();
5376 case Intrinsic::spv_undef: {
5377 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5383 case Intrinsic::spv_poison:
5384 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5389 case Intrinsic::spv_freeze:
5390 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5393 .
addUse(
I.getOperand(2).getReg())
5396 case Intrinsic::spv_named_boolean_spec_constant: {
5397 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5398 : SPIRV::OpSpecConstantFalse;
5400 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5401 .
addDef(
I.getOperand(0).getReg())
5404 unsigned SpecId =
I.getOperand(2).getImm();
5406 SPIRV::Decoration::SpecId, {SpecId});
5410 case Intrinsic::spv_const_composite: {
5412 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5418 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5420 std::function<bool(
Register)> HasSpecConstOperand =
5430 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5431 J < Def->getNumExplicitOperands(); ++J) {
5432 if (
Def->getOperand(J).isReg() &&
5433 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5439 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5440 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5441 : SPIRV::OpConstantComposite;
5442 unsigned ContinuedOpc = HasSpecConst
5443 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5444 : SPIRV::OpConstantCompositeContinuedINTEL;
5445 MachineIRBuilder MIR(
I);
5447 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5449 for (
auto *Instr : Instructions) {
5450 Instr->setDebugLoc(
I.getDebugLoc());
5455 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5462 case Intrinsic::spv_assign_name: {
5463 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5464 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5465 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5466 i <
I.getNumExplicitOperands(); ++i) {
5467 MIB.
addImm(
I.getOperand(i).getImm());
5472 case Intrinsic::spv_switch: {
5473 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5474 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5475 if (
I.getOperand(i).isReg())
5476 MIB.
addReg(
I.getOperand(i).getReg());
5477 else if (
I.getOperand(i).isCImm())
5478 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5479 else if (
I.getOperand(i).isMBB())
5480 MIB.
addMBB(
I.getOperand(i).getMBB());
5487 case Intrinsic::spv_loop_merge: {
5488 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5489 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5490 if (
I.getOperand(i).isMBB())
5491 MIB.
addMBB(
I.getOperand(i).getMBB());
5498 case Intrinsic::spv_loop_control_intel: {
5500 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5501 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5506 case Intrinsic::spv_selection_merge: {
5508 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5509 assert(
I.getOperand(1).isMBB() &&
5510 "operand 1 to spv_selection_merge must be a basic block");
5511 MIB.
addMBB(
I.getOperand(1).getMBB());
5512 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5516 case Intrinsic::spv_cmpxchg:
5517 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5518 case Intrinsic::spv_unreachable:
5519 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5522 case Intrinsic::spv_abort:
5523 return selectAbort(
I);
5524 case Intrinsic::spv_alloca:
5525 return selectFrameIndex(ResVReg, ResType,
I);
5526 case Intrinsic::spv_alloca_array:
5527 return selectAllocaArray(ResVReg, ResType,
I);
5528 case Intrinsic::spv_assume:
5530 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5531 .
addUse(
I.getOperand(1).getReg())
5536 case Intrinsic::spv_expect:
5538 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5541 .
addUse(
I.getOperand(2).getReg())
5542 .
addUse(
I.getOperand(3).getReg())
5547 case Intrinsic::arithmetic_fence:
5548 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5549 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5552 .
addUse(
I.getOperand(2).getReg())
5556 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5558 case Intrinsic::spv_thread_id:
5564 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5566 case Intrinsic::spv_thread_id_in_group:
5572 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5574 case Intrinsic::spv_group_id:
5580 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5582 case Intrinsic::spv_flattened_thread_id_in_group:
5589 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5591 case Intrinsic::spv_workgroup_size:
5592 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5594 case Intrinsic::spv_global_size:
5595 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5597 case Intrinsic::spv_global_offset:
5598 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5600 case Intrinsic::spv_num_workgroups:
5601 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5603 case Intrinsic::spv_subgroup_size:
5604 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5606 case Intrinsic::spv_num_subgroups:
5607 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5609 case Intrinsic::spv_subgroup_id:
5610 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5611 case Intrinsic::spv_subgroup_local_invocation_id:
5612 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5613 ResVReg, ResType,
I);
5614 case Intrinsic::spv_subgroup_max_size:
5615 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5617 case Intrinsic::spv_fdot:
5618 return selectFloatDot(ResVReg, ResType,
I);
5619 case Intrinsic::spv_udot:
5620 case Intrinsic::spv_sdot:
5621 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5623 return selectIntegerDot(ResVReg, ResType,
I,
5624 IID == Intrinsic::spv_sdot);
5625 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5626 case Intrinsic::spv_dot4add_i8packed:
5627 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5629 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5630 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5631 case Intrinsic::spv_dot4add_u8packed:
5632 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5634 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5635 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5636 case Intrinsic::spv_all:
5637 return selectAll(ResVReg, ResType,
I);
5638 case Intrinsic::spv_any:
5639 return selectAny(ResVReg, ResType,
I);
5640 case Intrinsic::spv_distance:
5641 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5642 case Intrinsic::spv_lerp:
5643 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5644 case Intrinsic::spv_length:
5645 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5646 case Intrinsic::spv_degrees:
5647 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5648 case Intrinsic::spv_faceforward:
5649 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5650 case Intrinsic::spv_frac:
5651 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5652 case Intrinsic::spv_isinf:
5653 return selectOpIsInf(ResVReg, ResType,
I);
5654 case Intrinsic::spv_isnan:
5655 return selectOpIsNan(ResVReg, ResType,
I);
5656 case Intrinsic::spv_isfinite:
5657 return selectOpIsFinite(ResVReg, ResType,
I);
5658 case Intrinsic::spv_isnormal:
5659 return selectOpIsNormal(ResVReg, ResType,
I);
5660 case Intrinsic::spv_normalize:
5661 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5662 case Intrinsic::spv_refract:
5663 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5664 case Intrinsic::spv_reflect:
5665 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5666 case Intrinsic::spv_rsqrt:
5667 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5668 case Intrinsic::spv_sign:
5669 return selectSign(ResVReg, ResType,
I);
5670 case Intrinsic::spv_smoothstep:
5671 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5672 case Intrinsic::spv_firstbituhigh:
5673 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5674 case Intrinsic::spv_firstbitshigh:
5675 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5676 case Intrinsic::spv_firstbitlow:
5677 return selectFirstBitLow(ResVReg, ResType,
I);
5678 case Intrinsic::spv_all_memory_barrier:
5679 return selectBarrierInst(
I, SPIRV::Scope::Device,
5680 SPIRV::MemorySemantics::UniformMemory |
5681 SPIRV::MemorySemantics::ImageMemory |
5682 SPIRV::MemorySemantics::WorkgroupMemory,
5684 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5685 return selectBarrierInst(
I, SPIRV::Scope::Device,
5686 SPIRV::MemorySemantics::UniformMemory |
5687 SPIRV::MemorySemantics::ImageMemory |
5688 SPIRV::MemorySemantics::WorkgroupMemory,
5690 case Intrinsic::spv_device_memory_barrier:
5691 return selectBarrierInst(
I, SPIRV::Scope::Device,
5692 SPIRV::MemorySemantics::UniformMemory |
5693 SPIRV::MemorySemantics::ImageMemory,
5695 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5696 return selectBarrierInst(
I, SPIRV::Scope::Device,
5697 SPIRV::MemorySemantics::UniformMemory |
5698 SPIRV::MemorySemantics::ImageMemory,
5700 case Intrinsic::spv_group_memory_barrier:
5701 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5702 SPIRV::MemorySemantics::WorkgroupMemory,
5704 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5705 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5706 SPIRV::MemorySemantics::WorkgroupMemory,
5708 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5709 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5710 SPIRV::StorageClass::StorageClass ResSC =
5713 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5714 "from the Generic storage class");
5715 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5723 case Intrinsic::spv_lifetime_start:
5724 case Intrinsic::spv_lifetime_end: {
5725 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5726 : SPIRV::OpLifetimeStop;
5727 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5728 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5737 case Intrinsic::spv_saturate:
5738 return selectSaturate(ResVReg, ResType,
I);
5739 case Intrinsic::spv_nclamp:
5740 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5741 case Intrinsic::spv_uclamp:
5742 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5743 case Intrinsic::spv_sclamp:
5744 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5745 case Intrinsic::spv_subgroup_prefix_bit_count:
5746 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5747 case Intrinsic::spv_wave_active_countbits:
5748 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5749 case Intrinsic::spv_wave_all_equal:
5750 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5751 case Intrinsic::spv_wave_all:
5752 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5753 case Intrinsic::spv_wave_any:
5754 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5755 case Intrinsic::spv_subgroup_ballot:
5756 return selectWaveOpInst(ResVReg, ResType,
I,
5757 SPIRV::OpGroupNonUniformBallot);
5758 case Intrinsic::spv_wave_is_first_lane:
5759 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5760 case Intrinsic::spv_wave_reduce_or:
5761 return selectWaveReduceOp(ResVReg, ResType,
I,
5762 SPIRV::OpGroupNonUniformBitwiseOr);
5763 case Intrinsic::spv_wave_reduce_xor:
5764 return selectWaveReduceOp(ResVReg, ResType,
I,
5765 SPIRV::OpGroupNonUniformBitwiseXor);
5766 case Intrinsic::spv_wave_reduce_and:
5767 return selectWaveReduceOp(ResVReg, ResType,
I,
5768 SPIRV::OpGroupNonUniformBitwiseAnd);
5769 case Intrinsic::spv_wave_reduce_umax:
5770 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5771 case Intrinsic::spv_wave_reduce_max:
5772 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5773 case Intrinsic::spv_wave_reduce_umin:
5774 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5775 case Intrinsic::spv_wave_reduce_min:
5776 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5777 case Intrinsic::spv_wave_reduce_sum:
5778 return selectWaveReduceSum(ResVReg, ResType,
I);
5779 case Intrinsic::spv_wave_product:
5780 return selectWaveReduceProduct(ResVReg, ResType,
I);
5781 case Intrinsic::spv_wave_readlane:
5782 return selectWaveOpInst(ResVReg, ResType,
I,
5783 SPIRV::OpGroupNonUniformShuffle);
5784 case Intrinsic::spv_wave_readlane_first:
5785 return selectWaveOpInst(ResVReg, ResType,
I,
5786 SPIRV::OpGroupNonUniformBroadcastFirst);
5787 case Intrinsic::spv_wave_prefix_sum:
5788 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5789 case Intrinsic::spv_wave_prefix_product:
5790 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5791 case Intrinsic::spv_quad_read_across_x: {
5792 return selectQuadSwap(ResVReg, ResType,
I, 0);
5794 case Intrinsic::spv_quad_read_across_y: {
5795 return selectQuadSwap(ResVReg, ResType,
I, 1);
5797 case Intrinsic::spv_quad_read_across_diagonal: {
5798 return selectQuadSwap(ResVReg, ResType,
I, 2);
5800 case Intrinsic::spv_radians:
5801 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5805 case Intrinsic::instrprof_increment:
5806 case Intrinsic::instrprof_increment_step:
5807 case Intrinsic::instrprof_value_profile:
5810 case Intrinsic::spv_value_md:
5812 case Intrinsic::spv_resource_handlefrombinding: {
5813 return selectHandleFromBinding(ResVReg, ResType,
I);
5815 case Intrinsic::spv_resource_counterhandlefrombinding:
5816 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5817 case Intrinsic::spv_resource_updatecounter:
5818 return selectUpdateCounter(ResVReg, ResType,
I);
5819 case Intrinsic::spv_resource_store_typedbuffer: {
5820 return selectImageWriteIntrinsic(
I);
5822 case Intrinsic::spv_resource_load_typedbuffer: {
5823 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5825 case Intrinsic::spv_resource_load_level: {
5826 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5828 case Intrinsic::spv_resource_getdimensions_x:
5829 case Intrinsic::spv_resource_getdimensions_xy:
5830 case Intrinsic::spv_resource_getdimensions_xyz: {
5831 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5833 case Intrinsic::spv_resource_getdimensions_levels_x:
5834 case Intrinsic::spv_resource_getdimensions_levels_xy:
5835 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5836 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5838 case Intrinsic::spv_resource_getdimensions_ms_xy:
5839 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5840 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5842 case Intrinsic::spv_resource_calculate_lod:
5843 case Intrinsic::spv_resource_calculate_lod_unclamped:
5844 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5845 case Intrinsic::spv_resource_sample:
5846 case Intrinsic::spv_resource_sample_clamp:
5847 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5848 case Intrinsic::spv_resource_samplebias:
5849 case Intrinsic::spv_resource_samplebias_clamp:
5850 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5851 case Intrinsic::spv_resource_samplegrad:
5852 case Intrinsic::spv_resource_samplegrad_clamp:
5853 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5854 case Intrinsic::spv_resource_samplelevel:
5855 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5856 case Intrinsic::spv_resource_samplecmp:
5857 case Intrinsic::spv_resource_samplecmp_clamp:
5858 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5859 case Intrinsic::spv_resource_samplecmplevelzero:
5860 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5861 case Intrinsic::spv_resource_gather:
5862 case Intrinsic::spv_resource_gather_cmp:
5863 return selectGatherIntrinsic(ResVReg, ResType,
I);
5864 case Intrinsic::spv_resource_getbasepointer:
5865 case Intrinsic::spv_resource_getpointer: {
5866 return selectResourceGetPointer(ResVReg, ResType,
I);
5868 case Intrinsic::spv_pushconstant_getpointer: {
5869 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5871 case Intrinsic::spv_discard: {
5872 return selectDiscard(ResVReg, ResType,
I);
5874 case Intrinsic::spv_resource_nonuniformindex: {
5875 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5877 case Intrinsic::spv_unpackhalf2x16: {
5878 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5880 case Intrinsic::spv_packhalf2x16: {
5881 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5883 case Intrinsic::spv_ddx:
5884 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5885 case Intrinsic::spv_ddy:
5886 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5887 case Intrinsic::spv_ddx_coarse:
5888 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5889 case Intrinsic::spv_ddy_coarse:
5890 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5891 case Intrinsic::spv_ddx_fine:
5892 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5893 case Intrinsic::spv_ddy_fine:
5894 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5895 case Intrinsic::spv_fwidth:
5896 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5897 case Intrinsic::spv_masked_gather:
5898 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5899 return selectMaskedGather(ResVReg, ResType,
I);
5900 return diagnoseUnsupported(
5901 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5902 case Intrinsic::spv_masked_scatter:
5903 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5904 return selectMaskedScatter(
I);
5905 return diagnoseUnsupported(
5906 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5907 case Intrinsic::returnaddress:
5908 case Intrinsic::frameaddress: {
5910 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5917 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5922bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5923 SPIRVTypeInst ResType,
5924 MachineInstr &
I)
const {
5927 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5934bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5935 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5937 assert(Intr.getIntrinsicID() ==
5938 Intrinsic::spv_resource_counterhandlefrombinding);
5941 Register MainHandleReg = Intr.getOperand(2).getReg();
5943 assert(MainHandleDef->getIntrinsicID() ==
5944 Intrinsic::spv_resource_handlefrombinding);
5948 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(),
MRI);
5949 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5950 std::string CounterName =
5955 MachineIRBuilder MIRBuilder(
I);
5957 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5959 ArraySize, IndexReg, CounterName, MIRBuilder);
5961 return BuildCOPY(ResVReg, CounterVarReg,
I);
5964bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5965 SPIRVTypeInst ResType,
5966 MachineInstr &
I)
const {
5968 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5970 Register CounterHandleReg = Intr.getOperand(2).getReg();
5971 int64_t IncrVal = Intr.getOperand(3).getImm();
5978 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5979 assert(CounterVarPointeeType &&
5980 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5981 "Counter variable must be a struct");
5983 SPIRV::StorageClass::StorageBuffer &&
5984 "Counter variable must be in the storage buffer storage class");
5986 "Counter variable must have exactly 1 member in the struct");
5987 const SPIRVTypeInst MemberType =
5990 "Counter variable struct must have a single i32 member");
5994 MachineIRBuilder MIRBuilder(
I);
5996 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5999 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
6005 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6008 .
addUse(CounterHandleReg)
6014 Register Scope = buildI32ConstantInEntryBlock(SPIRV::Scope::Device,
I);
6016 buildI32ConstantInEntryBlock(SPIRV::MemorySemantics::None,
I);
6019 buildI32ConstantInEntryBlock(
static_cast<uint32_t
>(IncrVal),
I);
6022 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
6031 return BuildCOPY(ResVReg, AtomicRes,
I);
6039 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
6047bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
6048 SPIRVTypeInst ResType,
6049 MachineInstr &
I)
const {
6057 Register ImageReg =
I.getOperand(2).getReg();
6065 Register IdxReg =
I.getOperand(3).getReg();
6067 MachineInstr &Pos =
I;
6069 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
6073bool SPIRVInstructionSelector::generateSampleImage(
6076 DebugLoc Loc, MachineInstr &Pos)
const {
6087 if (!loadHandleBeforePosition(NewSamplerReg,
6093 MachineIRBuilder MIRBuilder(Pos);
6106 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
6107 ImOps.Lod.has_value();
6108 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
6109 : SPIRV::OpImageSampleImplicitLod;
6111 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
6112 : SPIRV::OpImageSampleDrefImplicitLod;
6121 MIB.
addUse(*ImOps.Compare);
6123 uint32_t ImageOperands = 0;
6125 ImageOperands |= SPIRV::ImageOperand::Bias;
6127 ImageOperands |= SPIRV::ImageOperand::Lod;
6128 if (ImOps.GradX && ImOps.GradY)
6129 ImageOperands |= SPIRV::ImageOperand::Grad;
6130 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
6132 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6135 "Non-constant offsets are not supported in sample instructions.");
6140 ImageOperands |= SPIRV::ImageOperand::MinLod;
6142 if (ImageOperands != 0) {
6143 MIB.
addImm(ImageOperands);
6144 if (ImageOperands & SPIRV::ImageOperand::Bias)
6146 if (ImageOperands & SPIRV::ImageOperand::Lod)
6148 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6149 MIB.
addUse(*ImOps.GradX);
6150 MIB.
addUse(*ImOps.GradY);
6153 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6154 MIB.
addUse(*ImOps.Offset);
6155 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6156 MIB.
addUse(*ImOps.MinLod);
6163bool SPIRVInstructionSelector::selectImageQuerySize(
6165 std::optional<Register> LodReg)
const {
6168 "ImageReg is not an image type.");
6170 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6172 unsigned NumComponents = 0;
6174 case SPIRV::Dim::DIM_1D:
6175 case SPIRV::Dim::DIM_Buffer:
6176 NumComponents =
IsArray ? 2 : 1;
6178 case SPIRV::Dim::DIM_2D:
6179 case SPIRV::Dim::DIM_Cube:
6180 case SPIRV::Dim::DIM_Rect:
6181 NumComponents =
IsArray ? 3 : 2;
6183 case SPIRV::Dim::DIM_3D:
6187 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6197 if (!LodReg && IsSampled && !IsMultisampled &&
6198 Dim != SPIRV::Dim::DIM_Buffer && Dim != SPIRV::Dim::DIM_Rect)
6201 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6202 SPIRVTypeInst ResType =
6207 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6217bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6218 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6219 Register ImageReg =
I.getOperand(2).getReg();
6226 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6229bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6230 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6231 Register ImageReg =
I.getOperand(2).getReg();
6240 Register LodReg =
I.getOperand(3).getReg();
6243 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6245 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6252 TII.get(SPIRV::OpImageQueryLevels))
6259 TII.get(SPIRV::OpCompositeConstruct))
6269bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6270 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6271 Register ImageReg =
I.getOperand(2).getReg();
6282 "OpImageQuerySamples requires a multisampled image");
6284 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6292 TII.get(SPIRV::OpImageQuerySamples))
6299 TII.get(SPIRV::OpCompositeConstruct))
6309bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6310 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6311 Register ImageReg =
I.getOperand(2).getReg();
6312 Register SamplerReg =
I.getOperand(3).getReg();
6313 Register CoordinateReg =
I.getOperand(4).getReg();
6329 if (!loadHandleBeforePosition(
6334 MachineIRBuilder MIRBuilder(
I);
6340 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6350 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6357 unsigned ExtractedIndex =
6359 Intrinsic::spv_resource_calculate_lod_unclamped
6363 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6364 TII.get(SPIRV::OpCompositeExtract))
6374bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6375 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6376 Register ImageReg =
I.getOperand(2).getReg();
6377 Register SamplerReg =
I.getOperand(3).getReg();
6378 Register CoordinateReg =
I.getOperand(4).getReg();
6379 ImageOperands ImOps;
6380 if (
I.getNumOperands() > 5)
6381 ImOps.Offset =
I.getOperand(5).getReg();
6382 if (
I.getNumOperands() > 6)
6383 ImOps.MinLod =
I.getOperand(6).getReg();
6384 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6385 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6388bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6389 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6390 Register ImageReg =
I.getOperand(2).getReg();
6391 Register SamplerReg =
I.getOperand(3).getReg();
6392 Register CoordinateReg =
I.getOperand(4).getReg();
6393 ImageOperands ImOps;
6394 ImOps.Bias =
I.getOperand(5).getReg();
6395 if (
I.getNumOperands() > 6)
6396 ImOps.Offset =
I.getOperand(6).getReg();
6397 if (
I.getNumOperands() > 7)
6398 ImOps.MinLod =
I.getOperand(7).getReg();
6399 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6400 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6403bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6404 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6405 Register ImageReg =
I.getOperand(2).getReg();
6406 Register SamplerReg =
I.getOperand(3).getReg();
6407 Register CoordinateReg =
I.getOperand(4).getReg();
6408 ImageOperands ImOps;
6409 ImOps.GradX =
I.getOperand(5).getReg();
6410 ImOps.GradY =
I.getOperand(6).getReg();
6411 if (
I.getNumOperands() > 7)
6412 ImOps.Offset =
I.getOperand(7).getReg();
6413 if (
I.getNumOperands() > 8)
6414 ImOps.MinLod =
I.getOperand(8).getReg();
6415 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6416 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6419bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6420 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6421 Register ImageReg =
I.getOperand(2).getReg();
6422 Register SamplerReg =
I.getOperand(3).getReg();
6423 Register CoordinateReg =
I.getOperand(4).getReg();
6424 ImageOperands ImOps;
6425 ImOps.Lod =
I.getOperand(5).getReg();
6426 if (
I.getNumOperands() > 6)
6427 ImOps.Offset =
I.getOperand(6).getReg();
6428 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6429 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6432bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6433 SPIRVTypeInst ResType,
6434 MachineInstr &
I)
const {
6435 Register ImageReg =
I.getOperand(2).getReg();
6436 Register SamplerReg =
I.getOperand(3).getReg();
6437 Register CoordinateReg =
I.getOperand(4).getReg();
6438 ImageOperands ImOps;
6439 ImOps.Compare =
I.getOperand(5).getReg();
6440 if (
I.getNumOperands() > 6)
6441 ImOps.Offset =
I.getOperand(6).getReg();
6442 if (
I.getNumOperands() > 7)
6443 ImOps.MinLod =
I.getOperand(7).getReg();
6444 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6445 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6448bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6449 SPIRVTypeInst ResType,
6450 MachineInstr &
I)
const {
6451 Register ImageReg =
I.getOperand(2).getReg();
6452 Register CoordinateReg =
I.getOperand(3).getReg();
6453 Register LodReg =
I.getOperand(4).getReg();
6455 ImageOperands ImOps;
6457 if (
I.getNumOperands() > 5)
6458 ImOps.Offset =
I.getOperand(5).getReg();
6470 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6471 I.getDebugLoc(),
I, &ImOps);
6474bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6475 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6476 Register ImageReg =
I.getOperand(2).getReg();
6477 Register SamplerReg =
I.getOperand(3).getReg();
6478 Register CoordinateReg =
I.getOperand(4).getReg();
6479 ImageOperands ImOps;
6480 ImOps.Compare =
I.getOperand(5).getReg();
6481 if (
I.getNumOperands() > 6)
6482 ImOps.Offset =
I.getOperand(6).getReg();
6485 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6486 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6489bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6490 SPIRVTypeInst ResType,
6491 MachineInstr &
I)
const {
6492 Register ImageReg =
I.getOperand(2).getReg();
6493 Register SamplerReg =
I.getOperand(3).getReg();
6494 Register CoordinateReg =
I.getOperand(4).getReg();
6497 "ImageReg is not an image type.");
6502 ComponentOrCompareReg =
I.getOperand(5).getReg();
6503 OffsetReg =
I.getOperand(6).getReg();
6506 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6510 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6511 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6512 Dim != SPIRV::Dim::DIM_Rect) {
6514 "Gather operations are only supported for 2D, Cube, and Rect images.");
6521 if (!loadHandleBeforePosition(
6526 MachineIRBuilder MIRBuilder(
I);
6527 SPIRVTypeInst SampledImageType =
6532 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6540 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6542 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6544 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6549 .
addUse(ComponentOrCompareReg);
6551 uint32_t ImageOperands = 0;
6552 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6553 if (Dim == SPIRV::Dim::DIM_Cube) {
6555 "Gather operations with offset are not supported for Cube images.");
6559 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6561 ImageOperands |= SPIRV::ImageOperand::Offset;
6565 if (ImageOperands != 0) {
6566 MIB.
addImm(ImageOperands);
6568 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6576bool SPIRVInstructionSelector::generateImageReadOrFetch(
6579 const ImageOperands *ImOps)
const {
6582 "ImageReg is not an image type.");
6584 bool IsSignedInteger =
6589 bool IsFetch = (SampledOp.getImm() == 1);
6591 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6592 uint32_t ImageOperandsMask = 0;
6593 if (IsSignedInteger)
6594 ImageOperandsMask |= 0x1000;
6596 if (IsFetch && ImOps) {
6598 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6599 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6601 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6603 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6607 if (ImageOperandsMask != 0) {
6608 MIB.
addImm(ImageOperandsMask);
6609 if (IsFetch && ImOps) {
6612 if (ImOps->Offset &&
6613 (ImageOperandsMask &
6614 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6615 MIB.
addUse(*ImOps->Offset);
6624 SPIRVTypeInst SampledType =
6627 SPIRVTypeInst ReadType =
6628 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6629 bool ReadTypeMatchesResult = ReadType == ResType;
6631 Register ReadReg = ReadTypeMatchesResult
6637 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6643 BMI.constrainAllUses(
TII,
TRI, RBI);
6645 if (ReadTypeMatchesResult)
6658 if (ResultSize == 1) {
6667 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6670bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6671 SPIRVTypeInst ResType,
6672 MachineInstr &
I)
const {
6673 Register ResourcePtr =
I.getOperand(2).getReg();
6675 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6684 MachineIRBuilder MIRBuilder(
I);
6689 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6695 if (
I.getNumExplicitOperands() > 3) {
6696 Register IndexReg =
I.getOperand(3).getReg();
6703bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6704 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6709bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6710 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6711 Register ObjReg =
I.getOperand(2).getReg();
6712 if (!BuildCOPY(ResVReg, ObjReg,
I))
6722 decorateUsesAsNonUniform(ResVReg);
6726void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6729 {NonUniformReg,
nullptr}};
6730 llvm::SmallSet<Register, 8> Visited;
6731 while (WorkList.
size() > 0) {
6734 if (!Visited.
insert(CurrentReg).second)
6737 bool IsDecorated =
false;
6739 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6740 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6746 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6748 if (ResultReg == CurrentReg)
6756 MachineInstr &InsertPt =
6759 SPIRV::Decoration::NonUniformEXT, {});
6764bool SPIRVInstructionSelector::extractSubvector(
6766 MachineInstr &InsertionPoint)
const {
6768 [[maybe_unused]]
uint64_t InputSize =
6771 [[maybe_unused]]
bool IsLongVectorEXT =
6773 assert((InputSize > 1 || IsLongVectorEXT) &&
"The input must be a vector.");
6774 assert((ResultSize > 1 || IsLongVectorEXT) &&
"The result must be a vector.");
6775 assert(ResultSize < InputSize &&
6776 "Cannot extract more element than there are in the input.");
6783 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6792 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6794 TII.get(SPIRV::OpCompositeConstruct))
6798 for (
Register ComponentReg : ComponentRegisters)
6799 MIB.
addUse(ComponentReg);
6804bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6805 MachineInstr &
I)
const {
6812 Register ImageReg =
I.getOperand(1).getReg();
6820 Register CoordinateReg =
I.getOperand(2).getReg();
6821 Register DataReg =
I.getOperand(3).getReg();
6824 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6832Register SPIRVInstructionSelector::buildPointerToResource(
6833 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6834 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6835 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6837 if (ArraySize == 1) {
6838 SPIRVTypeInst PtrType =
6841 "SpirvResType did not have an explicit layout.");
6846 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6847 SPIRVTypeInst VarPointerType =
6850 VarPointerType, Set,
Binding, Name, MIRBuilder);
6852 SPIRVTypeInst ResPointerType =
6865bool SPIRVInstructionSelector::selectFirstBitSet16(
6866 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6867 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6869 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6873 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6876bool SPIRVInstructionSelector::selectFirstBitSet32(
6878 unsigned BitSetOpcode)
const {
6879 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6882 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6889bool SPIRVInstructionSelector::selectFirstBitSet64(
6891 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6905 if (ComponentCount > 2) {
6906 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6908 unsigned Opcode) ->
bool {
6909 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6913 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6917 MachineIRBuilder MIRBuilder(
I);
6919 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6923 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6929 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6939 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6940 SPIRV::OpVectorExtractDynamic))
6942 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6943 SPIRV::OpVectorExtractDynamic))
6947 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6948 TII.get(SPIRV::OpVectorShuffle))
6956 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6962 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6963 TII.get(SPIRV::OpVectorShuffle))
6971 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6991 SelectOp = SPIRV::OpSelectSISCond;
6992 AddOp = SPIRV::OpIAddS;
7000 SelectOp = SPIRV::OpSelectVIVCond;
7001 AddOp = SPIRV::OpIAddV;
7007 Register RegSecondaryOffset = Reg0;
7011 if (SwapPrimarySide) {
7012 PrimaryReg = LowReg;
7013 SecondaryReg = HighReg;
7014 RegPrimaryOffset = Reg0;
7015 RegSecondaryOffset = Reg32;
7020 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
7021 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
7026 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
7027 SPIRV::OpINotEqual))
7034 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
7035 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
7040 if (SwapPrimarySide) {
7042 if (!selectOpWithSrcs(RegAdd, ResType,
I,
7043 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
7054 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
7055 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
7060 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
7061 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
7064 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
7068bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
7069 SPIRVTypeInst ResType,
7071 bool IsSigned)
const {
7073 Register OpReg =
I.getOperand(2).getReg();
7076 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
7077 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
7081 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7083 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7085 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7088 return diagnoseUnsupported(
7090 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
7094bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
7095 SPIRVTypeInst ResType,
7096 MachineInstr &
I)
const {
7098 Register OpReg =
I.getOperand(2).getReg();
7103 unsigned ExtendOpcode = SPIRV::OpUConvert;
7104 unsigned BitSetOpcode = GL::FindILsb;
7108 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7110 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7112 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7115 return diagnoseUnsupported(
I,
7116 "spv_firstbitlow only supports 16,32,64 bits.");
7120bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
7121 SPIRVTypeInst ResType,
7122 MachineInstr &
I)
const {
7127 bool UseUntypedPointers =
7128 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7129 unsigned Opcode = UseUntypedPointers
7130 ? SPIRV::OpUntypedVariableLengthArrayINTEL
7131 : SPIRV::OpVariableLengthArrayINTEL;
7133 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
7139 if (UseUntypedPointers) {
7142 "untyped variable length array result must have a recorded element "
7147 MIB.
addUse(
I.getOperand(2).getReg());
7151 unsigned Alignment =
I.getOperand(3).getImm();
7165 while (!Worklist.
empty()) {
7167 switch (
T->getOpcode()) {
7168 case SPIRV::OpTypeInt:
7169 case SPIRV::OpTypeFloat:
7170 case SPIRV::OpTypePointer:
7172 case SPIRV::OpTypeVector:
7173 case SPIRV::OpTypeVectorIdEXT:
7174 case SPIRV::OpTypeMatrix:
7175 case SPIRV::OpTypeArray: {
7176 Register OperandReg =
T->getOperand(1).getReg();
7180 case SPIRV::OpTypeStruct:
7181 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7182 Register OperandReg =
T->getOperand(Idx).getReg();
7197 Register TypeReg = Ty->getOperand(0).getReg();
7198 if (!Visited.
insert(TypeReg).second)
7201 switch (Ty->getOpcode()) {
7202 case SPIRV::OpTypePointer:
7203 if (Ty->getOperand(1).getImm() == SPIRV::StorageClass::StorageBuffer)
7207 case SPIRV::OpTypeArray:
7208 case SPIRV::OpTypeRuntimeArray:
7211 case SPIRV::OpTypeStruct:
7212 for (
unsigned I = 1;
I < Ty->getNumOperands(); ++
I)
7228bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7229 assert(
I.getNumExplicitOperands() == 2);
7231 Register MsgReg =
I.getOperand(1).getReg();
7233 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7236 return diagnoseUnsupported(
7238 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7239 "scalar, pointer, vector, matrix, or aggregate of such types)");
7242 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7249bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7260 uint32_t MsgVal = ~0
u;
7261 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7262 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7265 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7268 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7275bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7276 SPIRVTypeInst ResType,
7277 MachineInstr &
I)
const {
7284 bool UseUntypedPointers =
7285 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7287 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7290 MachineIRBuilder MIRBuilder(
I);
7293 .
addImm(SPIRV::Extension::SPV_KHR_variable_pointers);
7295 .
addImm(SPIRV::Capability::VariablePointersStorageBuffer);
7298 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7301 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7305 if (UseUntypedPointers) {
7309 return diagnoseUnsupported(
7310 I,
"could not deduce the data type of an untyped variable");
7316 unsigned Alignment =
I.getOperand(2).getImm();
7323bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7328 const MachineInstr *PrevI =
I.getPrevNode();
7330 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7334 .
addMBB(
I.getOperand(0).getMBB())
7339 .
addMBB(
I.getOperand(0).getMBB())
7344bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7355 const MachineInstr *NextI =
I.getNextNode();
7357 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7363 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7365 .
addUse(
I.getOperand(0).getReg())
7366 .
addMBB(
I.getOperand(1).getMBB())
7372bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7373 MachineInstr &
I)
const {
7375 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7377 const unsigned NumOps =
I.getNumOperands();
7378 for (
unsigned i = 1; i <
NumOps; i += 2) {
7379 MIB.
addUse(
I.getOperand(i + 0).getReg());
7380 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7386bool SPIRVInstructionSelector::selectGlobalValue(
7387 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7389 MachineIRBuilder MIRBuilder(
I);
7390 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7393 std::string GlobalIdent;
7395 unsigned &
ID = UnnamedGlobalIDs[GV];
7397 ID = UnnamedGlobalIDs.
size();
7398 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7424 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7431 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
7436 MachineInstrBuilder MIB1 =
7437 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7440 MachineInstrBuilder MIB2 =
7442 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7446 GR.
add(ConstVal, MIB2);
7454 MachineInstrBuilder MIB3 =
7455 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7458 GR.
add(ConstVal, MIB3);
7464 assert(NewReg != ResVReg);
7465 return BuildCOPY(ResVReg, NewReg,
I);
7475 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7478 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7484 SPIRVTypeInst ResType =
7488 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7493 if (
GlobalVar->isExternallyInitialized() &&
7494 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7495 constexpr unsigned ReadWriteINTEL = 3u;
7498 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7504bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7505 SPIRVTypeInst ResType,
7506 MachineInstr &
I)
const {
7508 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7516 MachineIRBuilder MIRBuilder(
I);
7521 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7524 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7526 .
add(
I.getOperand(1))
7540 APFloat::rmNearestTiesToEven, &LosesInfo);
7545 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
7555bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7556 SPIRVTypeInst ResType,
7557 MachineInstr &
I)
const {
7560 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7566 Register ExpReg =
I.getOperand(2).getReg();
7568 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7569 SPIRV::OpConvertSToF))
7571 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7578bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7579 SPIRVTypeInst ResType,
7580 MachineInstr &
I)
const {
7596 MachineIRBuilder MIRBuilder(
I);
7597 SPIRVTypeInst FloatType =
7601 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7614 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7615 const bool IsUntyped =
7616 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7618 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7619 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7620 : SPIRV::OpVariable))
7623 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7631 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7634 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7637 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7641 Register IntegralPartReg =
I.getOperand(1).getReg();
7644 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7654 assert(
false &&
"GLSL::Modf is deprecated.");
7665bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7666 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7667 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7668 MachineIRBuilder MIRBuilder(
I);
7669 const SPIRVTypeInst Vec3Ty =
7672 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7684 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7688 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
7694 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7701 assert(
I.getOperand(2).isReg());
7702 const uint32_t ThreadId =
foldImm(
I.getOperand(2),
MRI);
7706 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7717bool SPIRVInstructionSelector::loadBuiltinInputID(
7718 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7719 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7720 MachineIRBuilder MIRBuilder(
I);
7722 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7737 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7741 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7750SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7751 MachineInstr &
I)
const {
7752 MachineIRBuilder MIRBuilder(
I);
7763bool SPIRVInstructionSelector::loadHandleBeforePosition(
7764 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7765 MachineInstr &Pos)
const {
7768 Intrinsic::spv_resource_handlefrombinding);
7776 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7777 MachineIRBuilder MIRBuilder(HandleDef);
7778 SPIRVTypeInst VarType = ResType;
7779 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7781 if (IsStructuredBuffer) {
7790 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7793 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7794 ArraySize, IndexReg, Name, MIRBuilder);
7798 uint32_t LoadOpcode =
7799 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7809bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7810 MachineInstr &
I)
const {
7812 return diagnoseUnsupported(
7813 I,
"this instruction is only supported in shaders.");
7818InstructionSelector *
7822 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