34#include "llvm/IR/IntrinsicsSPIRV.h"
39#define DEBUG_TYPE "spirv-isel"
46 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
51 std::optional<Register> Bias;
52 std::optional<Register>
Offset;
53 std::optional<Register> MinLod;
54 std::optional<Register> GradX;
55 std::optional<Register> GradY;
56 std::optional<Register> Lod;
57 std::optional<Register> Compare;
60llvm::SPIRV::SelectionControl::SelectionControl
61getSelectionOperandForImm(
int Imm) {
63 return SPIRV::SelectionControl::Flatten;
65 return SPIRV::SelectionControl::DontFlatten;
67 return SPIRV::SelectionControl::None;
71#define GET_GLOBALISEL_PREDICATE_BITSET
72#include "SPIRVGenGlobalISel.inc"
73#undef GET_GLOBALISEL_PREDICATE_BITSET
100#define GET_GLOBALISEL_PREDICATES_DECL
101#include "SPIRVGenGlobalISel.inc"
102#undef GET_GLOBALISEL_PREDICATES_DECL
104#define GET_GLOBALISEL_TEMPORARIES_DECL
105#include "SPIRVGenGlobalISel.inc"
106#undef GET_GLOBALISEL_TEMPORARIES_DECL
130 unsigned BitSetOpcode)
const;
134 unsigned BitSetOpcode)
const;
138 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
142 unsigned BitSetOpcode,
143 bool SwapPrimarySide)
const;
150 unsigned Opcode)
const;
153 unsigned Opcode)
const;
172 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
183 unsigned OpType)
const;
238 template <
bool Signed>
241 template <
bool Signed>
248 template <
typename PickOpcodeFn>
251 PickOpcodeFn &&PickOpcode)
const;
268 template <
typename PickOpcodeFn>
271 PickOpcodeFn &&PickOpcode)
const;
289 bool IsSigned)
const;
291 bool IsSigned,
unsigned Opcode)
const;
293 bool IsSigned)
const;
299 bool IsSigned)
const;
338 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
339 bool useMISrc =
true,
341 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
342 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
343 bool useMISrc =
true,
345 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
346 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
347 bool setMIFlags =
true,
bool useMISrc =
true,
349 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
350 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
351 bool useMISrc =
true,
354 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
355 MachineInstr &
I)
const;
357 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
358 MachineInstr &
I)
const;
360 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
361 MachineInstr &
I)
const;
363 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
364 MachineInstr &
I,
unsigned Opcode)
const;
366 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
367 bool WithGroupSync)
const;
369 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
370 MachineInstr &
I)
const;
372 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
373 MachineInstr &
I)
const;
377 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
378 MachineInstr &
I)
const;
380 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
381 MachineInstr &
I)
const;
383 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
384 MachineInstr &
I)
const;
385 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
386 MachineInstr &
I)
const;
387 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
388 SPIRVTypeInst ResType,
389 MachineInstr &
I)
const;
390 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
391 MachineInstr &
I)
const;
394 std::optional<Register> LodReg = std::nullopt)
const;
395 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
396 MachineInstr &
I)
const;
397 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
398 MachineInstr &
I)
const;
399 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
400 MachineInstr &
I)
const;
401 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
402 MachineInstr &
I)
const;
403 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
404 MachineInstr &
I)
const;
405 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
406 MachineInstr &
I)
const;
407 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
408 MachineInstr &
I)
const;
409 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
410 SPIRVTypeInst ResType,
411 MachineInstr &
I)
const;
412 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
414 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
415 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
416 MachineInstr &
I)
const;
417 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
418 MachineInstr &
I)
const;
419 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
420 MachineInstr &
I)
const;
421 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
422 MachineInstr &
I)
const;
423 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
424 MachineInstr &
I)
const;
425 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
426 MachineInstr &
I)
const;
427 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
428 MachineInstr &
I)
const;
429 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
430 MachineInstr &
I)
const;
431 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
432 MachineInstr &
I,
const unsigned DPdOpCode)
const;
434 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
435 SPIRVTypeInst ResType =
nullptr)
const;
437 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
438 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
439 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
441 MachineInstr &
I)
const;
442 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
444 bool wrapIntoSpecConstantOp(MachineInstr &
I,
447 Register getUcharPtrTypeReg(MachineInstr &
I,
448 SPIRV::StorageClass::StorageClass SC)
const;
449 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
451 uint32_t Opcode)
const;
452 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
453 SPIRVTypeInst SrcPtrTy)
const;
454 Register buildPointerToResource(SPIRVTypeInst ResType,
455 SPIRV::StorageClass::StorageClass SC,
456 uint32_t Set, uint32_t
Binding,
457 uint32_t ArraySize,
Register IndexReg,
459 MachineIRBuilder MIRBuilder)
const;
460 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
461 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
462 Register &ReadReg, MachineInstr &InsertionPoint)
const;
463 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
466 const ImageOperands *ImOps =
nullptr)
const;
467 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
469 Register CoordinateReg,
const ImageOperands &ImOps,
472 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
473 Register ResVReg, SPIRVTypeInst ResType,
474 MachineInstr &
I)
const;
475 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
476 Register ResVReg, SPIRVTypeInst ResType,
477 MachineInstr &
I)
const;
478 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
479 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
480 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
481 void errorIfInstrOutsideShader(MachineInstr &
I)
const;
484bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
486 if (
TET->getTargetExtName() ==
"spirv.Image") {
489 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
490 return TET->getTypeParameter(0)->isIntegerTy();
494#define GET_GLOBALISEL_IMPL
495#include "SPIRVGenGlobalISel.inc"
496#undef GET_GLOBALISEL_IMPL
502 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
505#include
"SPIRVGenGlobalISel.inc"
508#include
"SPIRVGenGlobalISel.inc"
520 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
524void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
525 if (HasVRegsReset == &MF)
540 for (
const auto &
MBB : MF) {
541 for (
const auto &
MI :
MBB) {
544 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
548 LLT DstType = MRI.
getType(DstReg);
550 LLT SrcType = MRI.
getType(SrcReg);
551 if (DstType != SrcType)
556 if (DstRC != SrcRC && SrcRC)
568 while (!Stack.empty()) {
573 switch (
MI->getOpcode()) {
574 case TargetOpcode::G_INTRINSIC:
575 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
576 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
579 if (IntrID != Intrinsic::spv_const_composite &&
580 IntrID != Intrinsic::spv_undef)
584 case TargetOpcode::G_BUILD_VECTOR:
585 case TargetOpcode::G_SPLAT_VECTOR:
587 i < OpDef->getNumOperands(); i++) {
592 Stack.push_back(OpNestedDef);
595 case TargetOpcode::G_CONSTANT:
596 case TargetOpcode::G_FCONSTANT:
597 case TargetOpcode::G_IMPLICIT_DEF:
598 case SPIRV::OpConstantTrue:
599 case SPIRV::OpConstantFalse:
600 case SPIRV::OpConstantI:
601 case SPIRV::OpConstantF:
602 case SPIRV::OpConstantComposite:
603 case SPIRV::OpConstantCompositeContinuedINTEL:
604 case SPIRV::OpConstantSampler:
605 case SPIRV::OpConstantNull:
607 case SPIRV::OpConstantFunctionPointerINTEL:
634 case Intrinsic::spv_all:
635 case Intrinsic::spv_alloca:
636 case Intrinsic::spv_any:
637 case Intrinsic::spv_bitcast:
638 case Intrinsic::spv_const_composite:
639 case Intrinsic::spv_cross:
640 case Intrinsic::spv_degrees:
641 case Intrinsic::spv_distance:
642 case Intrinsic::spv_extractelt:
643 case Intrinsic::spv_extractv:
644 case Intrinsic::spv_faceforward:
645 case Intrinsic::spv_fdot:
646 case Intrinsic::spv_firstbitlow:
647 case Intrinsic::spv_firstbitshigh:
648 case Intrinsic::spv_firstbituhigh:
649 case Intrinsic::spv_frac:
650 case Intrinsic::spv_gep:
651 case Intrinsic::spv_global_offset:
652 case Intrinsic::spv_global_size:
653 case Intrinsic::spv_group_id:
654 case Intrinsic::spv_insertelt:
655 case Intrinsic::spv_insertv:
656 case Intrinsic::spv_isinf:
657 case Intrinsic::spv_isnan:
658 case Intrinsic::spv_lerp:
659 case Intrinsic::spv_length:
660 case Intrinsic::spv_normalize:
661 case Intrinsic::spv_num_subgroups:
662 case Intrinsic::spv_num_workgroups:
663 case Intrinsic::spv_ptrcast:
664 case Intrinsic::spv_radians:
665 case Intrinsic::spv_reflect:
666 case Intrinsic::spv_refract:
667 case Intrinsic::spv_resource_getpointer:
668 case Intrinsic::spv_resource_handlefrombinding:
669 case Intrinsic::spv_resource_handlefromimplicitbinding:
670 case Intrinsic::spv_resource_nonuniformindex:
671 case Intrinsic::spv_resource_sample:
672 case Intrinsic::spv_rsqrt:
673 case Intrinsic::spv_saturate:
674 case Intrinsic::spv_sdot:
675 case Intrinsic::spv_sign:
676 case Intrinsic::spv_smoothstep:
677 case Intrinsic::spv_step:
678 case Intrinsic::spv_subgroup_id:
679 case Intrinsic::spv_subgroup_local_invocation_id:
680 case Intrinsic::spv_subgroup_max_size:
681 case Intrinsic::spv_subgroup_size:
682 case Intrinsic::spv_thread_id:
683 case Intrinsic::spv_thread_id_in_group:
684 case Intrinsic::spv_udot:
685 case Intrinsic::spv_undef:
686 case Intrinsic::spv_value_md:
687 case Intrinsic::spv_workgroup_size:
699 case SPIRV::OpTypeVoid:
700 case SPIRV::OpTypeBool:
701 case SPIRV::OpTypeInt:
702 case SPIRV::OpTypeFloat:
703 case SPIRV::OpTypeVector:
704 case SPIRV::OpTypeMatrix:
705 case SPIRV::OpTypeImage:
706 case SPIRV::OpTypeSampler:
707 case SPIRV::OpTypeSampledImage:
708 case SPIRV::OpTypeArray:
709 case SPIRV::OpTypeRuntimeArray:
710 case SPIRV::OpTypeStruct:
711 case SPIRV::OpTypeOpaque:
712 case SPIRV::OpTypePointer:
713 case SPIRV::OpTypeFunction:
714 case SPIRV::OpTypeEvent:
715 case SPIRV::OpTypeDeviceEvent:
716 case SPIRV::OpTypeReserveId:
717 case SPIRV::OpTypeQueue:
718 case SPIRV::OpTypePipe:
719 case SPIRV::OpTypeForwardPointer:
720 case SPIRV::OpTypePipeStorage:
721 case SPIRV::OpTypeNamedBarrier:
722 case SPIRV::OpTypeAccelerationStructureNV:
723 case SPIRV::OpTypeCooperativeMatrixNV:
724 case SPIRV::OpTypeCooperativeMatrixKHR:
734 if (
MI.getNumDefs() == 0)
737 for (
const auto &MO :
MI.all_defs()) {
739 if (
Reg.isPhysical()) {
744 if (
UseMI.getOpcode() != SPIRV::OpName) {
751 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
752 MI.isLifetimeMarker()) {
755 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
766 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
767 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
770 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
775 if (
MI.mayStore() ||
MI.isCall() ||
776 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
777 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
778 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
789 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
796void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
798 for (
const auto &MO :
MI.all_defs()) {
802 SmallVector<MachineInstr *, 4> UselessOpNames;
805 "There is still a use of the dead function.");
808 for (MachineInstr *OpNameMI : UselessOpNames) {
810 OpNameMI->eraseFromParent();
815void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
818 removeOpNamesForDeadMI(
MI);
819 MI.eraseFromParent();
822bool SPIRVInstructionSelector::select(MachineInstr &
I) {
823 resetVRegsType(*
I.getParent()->getParent());
825 assert(
I.getParent() &&
"Instruction should be in a basic block!");
826 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
831 removeDeadInstruction(
I);
838 if (Opcode == SPIRV::ASSIGN_TYPE) {
839 Register DstReg =
I.getOperand(0).getReg();
840 Register SrcReg =
I.getOperand(1).getReg();
843 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
844 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
845 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
846 Register SelectDstReg =
Def->getOperand(0).getReg();
847 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
849 assert(SuccessToSelectSelect);
851 Def->eraseFromParent();
858 bool Res = selectImpl(
I, *CoverageInfo);
860 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
861 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
865 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
877 }
else if (
I.getNumDefs() == 1) {
889 removeDeadInstruction(
I);
894 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
895 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
901 bool HasDefs =
I.getNumDefs() > 0;
904 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
905 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
906 if (spvSelect(ResVReg, ResType,
I)) {
908 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
919 case TargetOpcode::G_CONSTANT:
920 case TargetOpcode::G_FCONSTANT:
927 MachineInstr &
I)
const {
930 if (DstRC != SrcRC && SrcRC)
932 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
939bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
940 SPIRVTypeInst ResType,
941 MachineInstr &
I)
const {
942 const unsigned Opcode =
I.getOpcode();
944 return selectImpl(
I, *CoverageInfo);
946 case TargetOpcode::G_CONSTANT:
947 case TargetOpcode::G_FCONSTANT:
948 return selectConst(ResVReg, ResType,
I);
949 case TargetOpcode::G_GLOBAL_VALUE:
950 return selectGlobalValue(ResVReg,
I);
951 case TargetOpcode::G_IMPLICIT_DEF:
952 return selectOpUndef(ResVReg, ResType,
I);
953 case TargetOpcode::G_FREEZE:
954 return selectFreeze(ResVReg, ResType,
I);
956 case TargetOpcode::G_INTRINSIC:
957 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
958 case TargetOpcode::G_INTRINSIC_CONVERGENT:
959 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
960 return selectIntrinsic(ResVReg, ResType,
I);
961 case TargetOpcode::G_BITREVERSE:
962 return selectBitreverse(ResVReg, ResType,
I);
964 case TargetOpcode::G_BUILD_VECTOR:
965 return selectBuildVector(ResVReg, ResType,
I);
966 case TargetOpcode::G_SPLAT_VECTOR:
967 return selectSplatVector(ResVReg, ResType,
I);
969 case TargetOpcode::G_SHUFFLE_VECTOR: {
970 MachineBasicBlock &BB = *
I.getParent();
971 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
974 .
addUse(
I.getOperand(1).getReg())
975 .
addUse(
I.getOperand(2).getReg());
976 for (
auto V :
I.getOperand(3).getShuffleMask())
981 case TargetOpcode::G_MEMMOVE:
982 case TargetOpcode::G_MEMCPY:
983 case TargetOpcode::G_MEMSET:
984 return selectMemOperation(ResVReg,
I);
986 case TargetOpcode::G_ICMP:
987 return selectICmp(ResVReg, ResType,
I);
988 case TargetOpcode::G_FCMP:
989 return selectFCmp(ResVReg, ResType,
I);
991 case TargetOpcode::G_FRAME_INDEX:
992 return selectFrameIndex(ResVReg, ResType,
I);
994 case TargetOpcode::G_LOAD:
995 return selectLoad(ResVReg, ResType,
I);
996 case TargetOpcode::G_STORE:
997 return selectStore(
I);
999 case TargetOpcode::G_BR:
1000 return selectBranch(
I);
1001 case TargetOpcode::G_BRCOND:
1002 return selectBranchCond(
I);
1004 case TargetOpcode::G_PHI:
1005 return selectPhi(ResVReg,
I);
1007 case TargetOpcode::G_FPTOSI:
1008 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1009 case TargetOpcode::G_FPTOUI:
1010 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1012 case TargetOpcode::G_FPTOSI_SAT:
1013 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1014 case TargetOpcode::G_FPTOUI_SAT:
1015 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1017 case TargetOpcode::G_SITOFP:
1018 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1019 case TargetOpcode::G_UITOFP:
1020 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1022 case TargetOpcode::G_CTPOP:
1023 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitCount);
1024 case TargetOpcode::G_SMIN:
1025 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1026 case TargetOpcode::G_UMIN:
1027 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1029 case TargetOpcode::G_SMAX:
1030 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1031 case TargetOpcode::G_UMAX:
1032 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1034 case TargetOpcode::G_SCMP:
1035 return selectSUCmp(ResVReg, ResType,
I,
true);
1036 case TargetOpcode::G_UCMP:
1037 return selectSUCmp(ResVReg, ResType,
I,
false);
1038 case TargetOpcode::G_LROUND:
1039 case TargetOpcode::G_LLROUND: {
1042 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1044 regForLround, *(
I.getParent()->getParent()));
1046 CL::round, GL::Round,
false);
1048 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1055 case TargetOpcode::G_STRICT_FMA:
1056 case TargetOpcode::G_FMA: {
1059 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1062 .
addUse(
I.getOperand(1).getReg())
1063 .
addUse(
I.getOperand(2).getReg())
1064 .
addUse(
I.getOperand(3).getReg())
1069 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1072 case TargetOpcode::G_STRICT_FLDEXP:
1073 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1075 case TargetOpcode::G_FPOW:
1076 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1077 case TargetOpcode::G_FPOWI:
1078 return selectFpowi(ResVReg, ResType,
I);
1080 case TargetOpcode::G_FEXP:
1081 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1082 case TargetOpcode::G_FEXP2:
1083 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1084 case TargetOpcode::G_FEXP10:
1085 return selectExp10(ResVReg, ResType,
I);
1087 case TargetOpcode::G_FMODF:
1088 return selectModf(ResVReg, ResType,
I);
1089 case TargetOpcode::G_FSINCOS:
1090 return selectSincos(ResVReg, ResType,
I);
1092 case TargetOpcode::G_FLOG:
1093 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1094 case TargetOpcode::G_FLOG2:
1095 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1096 case TargetOpcode::G_FLOG10:
1097 return selectLog10(ResVReg, ResType,
I);
1099 case TargetOpcode::G_FABS:
1100 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1101 case TargetOpcode::G_ABS:
1102 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1104 case TargetOpcode::G_FMINNUM:
1105 case TargetOpcode::G_FMINIMUM:
1106 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1107 case TargetOpcode::G_FMAXNUM:
1108 case TargetOpcode::G_FMAXIMUM:
1109 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1111 case TargetOpcode::G_FCOPYSIGN:
1112 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1114 case TargetOpcode::G_FCEIL:
1115 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1116 case TargetOpcode::G_FFLOOR:
1117 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1119 case TargetOpcode::G_FCOS:
1120 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1121 case TargetOpcode::G_FSIN:
1122 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1123 case TargetOpcode::G_FTAN:
1124 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1125 case TargetOpcode::G_FACOS:
1126 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1127 case TargetOpcode::G_FASIN:
1128 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1129 case TargetOpcode::G_FATAN:
1130 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1131 case TargetOpcode::G_FATAN2:
1132 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1133 case TargetOpcode::G_FCOSH:
1134 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1135 case TargetOpcode::G_FSINH:
1136 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1137 case TargetOpcode::G_FTANH:
1138 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1140 case TargetOpcode::G_STRICT_FSQRT:
1141 case TargetOpcode::G_FSQRT:
1142 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1144 case TargetOpcode::G_CTTZ:
1145 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1146 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1147 case TargetOpcode::G_CTLZ:
1148 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1149 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1151 case TargetOpcode::G_INTRINSIC_ROUND:
1152 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1153 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1154 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1155 case TargetOpcode::G_INTRINSIC_TRUNC:
1156 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1157 case TargetOpcode::G_FRINT:
1158 case TargetOpcode::G_FNEARBYINT:
1159 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1161 case TargetOpcode::G_SMULH:
1162 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1163 case TargetOpcode::G_UMULH:
1164 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1166 case TargetOpcode::G_SADDSAT:
1167 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1168 case TargetOpcode::G_UADDSAT:
1169 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1170 case TargetOpcode::G_SSUBSAT:
1171 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1172 case TargetOpcode::G_USUBSAT:
1173 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1175 case TargetOpcode::G_FFREXP:
1176 return selectFrexp(ResVReg, ResType,
I);
1178 case TargetOpcode::G_UADDO:
1179 return selectOverflowArith(ResVReg, ResType,
I,
1180 ResType->
getOpcode() == SPIRV::OpTypeVector
1181 ? SPIRV::OpIAddCarryV
1182 : SPIRV::OpIAddCarryS);
1183 case TargetOpcode::G_USUBO:
1184 return selectOverflowArith(ResVReg, ResType,
I,
1185 ResType->
getOpcode() == SPIRV::OpTypeVector
1186 ? SPIRV::OpISubBorrowV
1187 : SPIRV::OpISubBorrowS);
1188 case TargetOpcode::G_UMULO:
1189 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1190 case TargetOpcode::G_SMULO:
1191 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1193 case TargetOpcode::G_SEXT:
1194 return selectExt(ResVReg, ResType,
I,
true);
1195 case TargetOpcode::G_ANYEXT:
1196 case TargetOpcode::G_ZEXT:
1197 return selectExt(ResVReg, ResType,
I,
false);
1198 case TargetOpcode::G_TRUNC:
1199 return selectTrunc(ResVReg, ResType,
I);
1200 case TargetOpcode::G_FPTRUNC:
1201 case TargetOpcode::G_FPEXT:
1202 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1204 case TargetOpcode::G_PTRTOINT:
1205 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1206 case TargetOpcode::G_INTTOPTR:
1207 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1208 case TargetOpcode::G_BITCAST:
1209 return selectBitcast(ResVReg, ResType,
I);
1210 case TargetOpcode::G_ADDRSPACE_CAST:
1211 return selectAddrSpaceCast(ResVReg, ResType,
I);
1212 case TargetOpcode::G_PTR_ADD: {
1214 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1218 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1219 (*II).getOpcode() == TargetOpcode::COPY ||
1220 (*II).getOpcode() == SPIRV::OpVariable) &&
1221 getImm(
I.getOperand(2), MRI));
1223 bool IsGVInit =
false;
1227 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1228 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1229 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1230 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1240 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1253 "incompatible result and operand types in a bitcast");
1255 MachineInstrBuilder MIB =
1256 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1263 : SPIRV::OpInBoundsPtrAccessChain))
1267 .
addUse(
I.getOperand(2).getReg())
1270 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1274 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1276 .
addUse(
I.getOperand(2).getReg())
1285 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1288 .
addImm(
static_cast<uint32_t
>(
1289 SPIRV::Opcode::InBoundsPtrAccessChain))
1292 .
addUse(
I.getOperand(2).getReg());
1297 case TargetOpcode::G_ATOMICRMW_OR:
1298 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1299 case TargetOpcode::G_ATOMICRMW_ADD:
1300 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1301 case TargetOpcode::G_ATOMICRMW_AND:
1302 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1303 case TargetOpcode::G_ATOMICRMW_MAX:
1304 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1305 case TargetOpcode::G_ATOMICRMW_MIN:
1306 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1307 case TargetOpcode::G_ATOMICRMW_SUB:
1308 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1309 case TargetOpcode::G_ATOMICRMW_XOR:
1310 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1311 case TargetOpcode::G_ATOMICRMW_UMAX:
1312 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1313 case TargetOpcode::G_ATOMICRMW_UMIN:
1314 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1315 case TargetOpcode::G_ATOMICRMW_XCHG:
1316 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1317 case TargetOpcode::G_ATOMIC_CMPXCHG:
1318 return selectAtomicCmpXchg(ResVReg, ResType,
I);
1320 case TargetOpcode::G_ATOMICRMW_FADD:
1321 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1322 case TargetOpcode::G_ATOMICRMW_FSUB:
1324 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1325 ResType->
getOpcode() == SPIRV::OpTypeVector
1327 : SPIRV::OpFNegate);
1328 case TargetOpcode::G_ATOMICRMW_FMIN:
1329 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1330 case TargetOpcode::G_ATOMICRMW_FMAX:
1331 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1333 case TargetOpcode::G_FENCE:
1334 return selectFence(
I);
1336 case TargetOpcode::G_STACKSAVE:
1337 return selectStackSave(ResVReg, ResType,
I);
1338 case TargetOpcode::G_STACKRESTORE:
1339 return selectStackRestore(
I);
1341 case TargetOpcode::G_UNMERGE_VALUES:
1347 case TargetOpcode::G_TRAP:
1348 case TargetOpcode::G_UBSANTRAP:
1349 case TargetOpcode::DBG_LABEL:
1351 case TargetOpcode::G_DEBUGTRAP:
1352 return selectDebugTrap(ResVReg, ResType,
I);
1359bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1360 SPIRVTypeInst ResType,
1361 MachineInstr &
I)
const {
1362 unsigned Opcode = SPIRV::OpNop;
1369bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1370 SPIRVTypeInst ResType,
1372 GL::GLSLExtInst GLInst,
1373 bool setMIFlags,
bool useMISrc,
1376 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1377 std::string DiagMsg;
1378 raw_string_ostream OS(DiagMsg);
1379 I.print(OS,
true,
false,
false,
false);
1380 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1383 return selectExtInst(ResVReg, ResType,
I,
1384 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1385 setMIFlags, useMISrc, SrcRegs);
1388bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1389 SPIRVTypeInst ResType,
1391 CL::OpenCLExtInst CLInst,
1392 bool setMIFlags,
bool useMISrc,
1394 return selectExtInst(ResVReg, ResType,
I,
1395 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1396 setMIFlags, useMISrc, SrcRegs);
1399bool SPIRVInstructionSelector::selectExtInst(
1400 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1401 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1403 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1404 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1405 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1409bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1410 SPIRVTypeInst ResType,
1413 bool setMIFlags,
bool useMISrc,
1416 for (
const auto &[InstructionSet, Opcode] : Insts) {
1420 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1423 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1428 const unsigned NumOps =
I.getNumOperands();
1431 I.getOperand(Index).getType() ==
1432 MachineOperand::MachineOperandType::MO_IntrinsicID)
1435 MIB.
add(
I.getOperand(Index));
1447bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1448 SPIRVTypeInst ResType,
1449 MachineInstr &
I)
const {
1450 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1451 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1452 for (
const auto &Ex : ExtInsts) {
1453 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1454 uint32_t Opcode = Ex.second;
1458 MachineIRBuilder MIRBuilder(
I);
1461 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1466 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1469 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1472 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1475 .
addImm(
static_cast<uint32_t
>(Ex.first))
1477 .
add(
I.getOperand(2))
1481 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1482 .
addDef(
I.getOperand(1).getReg())
1491bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1492 SPIRVTypeInst ResType,
1493 MachineInstr &
I)
const {
1494 Register CosResVReg =
I.getOperand(1).getReg();
1495 unsigned SrcIdx =
I.getNumExplicitDefs();
1500 MachineIRBuilder MIRBuilder(
I);
1502 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1507 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1510 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1512 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1515 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1517 .
add(
I.getOperand(SrcIdx))
1520 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1528 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1531 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1533 .
add(
I.getOperand(SrcIdx))
1535 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1538 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1540 .
add(
I.getOperand(SrcIdx))
1547bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1548 SPIRVTypeInst ResType,
1550 std::vector<Register> Srcs,
1551 unsigned Opcode)
const {
1552 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1562bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1563 SPIRVTypeInst ResType,
1565 unsigned Opcode)
const {
1567 Register SrcReg =
I.getOperand(1).getReg();
1572 unsigned DefOpCode = DefIt->getOpcode();
1573 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1576 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1577 DefOpCode = VRD->getOpcode();
1579 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1580 DefOpCode == TargetOpcode::G_CONSTANT ||
1581 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1587 uint32_t SpecOpcode = 0;
1589 case SPIRV::OpConvertPtrToU:
1590 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1592 case SPIRV::OpConvertUToPtr:
1593 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1598 TII.get(SPIRV::OpSpecConstantOp))
1608 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1612bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1613 SPIRVTypeInst ResType,
1614 MachineInstr &
I)
const {
1615 Register OpReg =
I.getOperand(1).getReg();
1616 SPIRVTypeInst OpType =
1620 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1630 if (
MemOp->isVolatile())
1631 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1632 if (
MemOp->isNonTemporal())
1633 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1635 if (!ST->isShader() &&
MemOp->getAlign().value())
1636 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1640 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1641 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1645 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1647 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1651 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1655 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1657 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1669 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1671 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1673 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1677bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1678 SPIRVTypeInst ResType,
1679 MachineInstr &
I)
const {
1681 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1686 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1687 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1689 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1693 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1697 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1698 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1699 I.getDebugLoc(),
I);
1703 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1707 if (!
I.getNumMemOperands()) {
1708 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1710 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1713 MachineIRBuilder MIRBuilder(
I);
1720bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1722 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1723 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1728 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1729 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1734 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1738 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1739 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1740 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1741 TII.get(SPIRV::OpImageWrite))
1747 if (sampledTypeIsSignedInteger(LLVMHandleType))
1750 BMI.constrainAllUses(
TII,
TRI, RBI);
1756 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpStore))
1759 if (!
I.getNumMemOperands()) {
1760 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1762 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1765 MachineIRBuilder MIRBuilder(
I);
1772bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
1773 SPIRVTypeInst ResType,
1774 MachineInstr &
I)
const {
1775 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
1783 const Register PtrsReg =
I.getOperand(2).getReg();
1784 const uint32_t Alignment =
I.getOperand(3).getImm();
1785 const Register MaskReg =
I.getOperand(4).getReg();
1786 const Register PassthruReg =
I.getOperand(5).getReg();
1787 const Register AlignmentReg = buildI32Constant(Alignment,
I);
1791 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
1802bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
1803 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
1810 const Register ValuesReg =
I.getOperand(1).getReg();
1811 const Register PtrsReg =
I.getOperand(2).getReg();
1812 const uint32_t Alignment =
I.getOperand(3).getImm();
1813 const Register MaskReg =
I.getOperand(4).getReg();
1814 const Register AlignmentReg = buildI32Constant(Alignment,
I);
1818 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
1827bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
1828 const Twine &Msg)
const {
1829 const Function &
F =
I.getMF()->getFunction();
1830 F.getContext().diagnose(
1831 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
1835bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
1836 SPIRVTypeInst ResType,
1837 MachineInstr &
I)
const {
1838 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1840 "llvm.stacksave intrinsic: this instruction requires the following "
1841 "SPIR-V extension: SPV_INTEL_variable_length_array",
1844 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
1851bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
1852 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1854 "llvm.stackrestore intrinsic: this instruction requires the following "
1855 "SPIR-V extension: SPV_INTEL_variable_length_array",
1857 if (!
I.getOperand(0).isReg())
1860 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
1861 .
addUse(
I.getOperand(0).getReg())
1867SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
1868 MachineIRBuilder MIRBuilder(
I);
1869 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1876 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
1880 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1881 Type *ArrTy = ArrayType::get(ValTy, Num);
1883 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
1886 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
1893 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
1896 .
addImm(SPIRV::StorageClass::UniformConstant)
1907bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
1910 Register DstReg =
I.getOperand(0).getReg();
1915 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
1920 "Unable to determine pointee type size for OpCopyMemory");
1921 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
1922 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
1924 "OpCopyMemory requires the size to match the pointee type size");
1925 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
1928 if (
I.getNumMemOperands()) {
1929 MachineIRBuilder MIRBuilder(
I);
1936bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
1939 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
1940 .
addUse(
I.getOperand(0).getReg())
1942 .
addUse(
I.getOperand(2).getReg());
1943 if (
I.getNumMemOperands()) {
1944 MachineIRBuilder MIRBuilder(
I);
1951bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
1952 MachineInstr &
I)
const {
1953 Register SrcReg =
I.getOperand(1).getReg();
1954 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
1955 Register VarReg = getOrCreateMemSetGlobal(
I);
1958 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1960 ValTy,
I, SPIRV::StorageClass::UniformConstant);
1962 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
1966 if (!selectCopyMemory(
I, SrcReg))
1969 if (!selectCopyMemorySized(
I, SrcReg))
1972 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
1973 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
1978bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
1979 SPIRVTypeInst ResType,
1982 unsigned NegateOpcode)
const {
1984 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1987 Register ScopeReg = buildI32Constant(Scope,
I);
1989 Register Ptr =
I.getOperand(1).getReg();
1995 Register MemSemReg = buildI32Constant(MemSem ,
I);
1997 Register ValueReg =
I.getOperand(2).getReg();
1998 if (NegateOpcode != 0) {
2001 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2006 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2017bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2018 unsigned ArgI =
I.getNumOperands() - 1;
2020 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2021 SPIRVTypeInst SrcType =
2023 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2025 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2027 SPIRVTypeInst ScalarType =
2030 unsigned CurrentIndex = 0;
2031 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2032 Register ResVReg =
I.getOperand(i).getReg();
2035 LLT ResLLT = MRI->
getType(ResVReg);
2041 ResType = ScalarType;
2047 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2050 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2056 for (
unsigned j = 0;
j < NumElements; ++
j) {
2057 MIB.
addImm(CurrentIndex + j);
2059 CurrentIndex += NumElements;
2063 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2075bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2078 Register MemSemReg = buildI32Constant(MemSem,
I);
2080 uint32_t
Scope =
static_cast<uint32_t
>(
2082 Register ScopeReg = buildI32Constant(Scope,
I);
2084 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2091bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2092 SPIRVTypeInst ResType,
2094 unsigned Opcode)
const {
2095 Type *ResTy =
nullptr;
2099 "Not enough info to select the arithmetic with overflow instruction");
2102 "with overflow instruction");
2108 MachineIRBuilder MIRBuilder(
I);
2110 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2111 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2117 Register ZeroReg = buildZerosVal(ResType,
I);
2122 if (ResName.
size() > 0)
2127 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2130 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2131 MIB.
addUse(
I.getOperand(i).getReg());
2136 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2137 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2139 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2140 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2147 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2148 .
addDef(
I.getOperand(1).getReg())
2156bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2157 SPIRVTypeInst ResType,
2158 MachineInstr &
I)
const {
2162 Register Ptr =
I.getOperand(2).getReg();
2165 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2168 ScopeReg = buildI32Constant(Scope,
I);
2170 unsigned ScSem =
static_cast<uint32_t
>(
2173 unsigned MemSemEq =
static_cast<uint32_t
>(
getMemSemantics(AO)) | ScSem;
2174 Register MemSemEqReg = buildI32Constant(MemSemEq,
I);
2176 unsigned MemSemNeq =
static_cast<uint32_t
>(
getMemSemantics(FO)) | ScSem;
2177 if (MemSemEq == MemSemNeq)
2178 MemSemNeqReg = MemSemEqReg;
2180 MemSemNeqReg = buildI32Constant(MemSemEq,
I);
2183 ScopeReg =
I.getOperand(5).getReg();
2184 MemSemEqReg =
I.getOperand(6).getReg();
2185 MemSemNeqReg =
I.getOperand(7).getReg();
2189 Register Val =
I.getOperand(4).getReg();
2193 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2212 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2219 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2231 case SPIRV::StorageClass::DeviceOnlyINTEL:
2232 case SPIRV::StorageClass::HostOnlyINTEL:
2241 bool IsGRef =
false;
2242 bool IsAllowedRefs =
2244 unsigned Opcode = It.getOpcode();
2245 if (Opcode == SPIRV::OpConstantComposite ||
2246 Opcode == SPIRV::OpSpecConstantComposite ||
2247 Opcode == SPIRV::OpVariable ||
2248 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2249 return IsGRef = true;
2250 return Opcode == SPIRV::OpName;
2252 return IsAllowedRefs && IsGRef;
2255Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2256 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2258 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2262SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2264 uint32_t Opcode)
const {
2265 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2266 TII.get(SPIRV::OpSpecConstantOp))
2274SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2275 SPIRVTypeInst SrcPtrTy)
const {
2276 SPIRVTypeInst GenericPtrTy =
2280 SPIRV::StorageClass::Generic),
2282 MachineFunction *MF =
I.getParent()->getParent();
2284 MachineInstrBuilder MIB = buildSpecConstantOp(
2286 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2296bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2297 SPIRVTypeInst ResType,
2298 MachineInstr &
I)
const {
2302 Register SrcPtr =
I.getOperand(1).getReg();
2306 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2307 ResType->
getOpcode() != SPIRV::OpTypePointer)
2308 return BuildCOPY(ResVReg, SrcPtr,
I);
2318 unsigned SpecOpcode =
2320 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2323 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2330 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2332 .constrainAllUses(
TII,
TRI, RBI);
2334 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2336 buildSpecConstantOp(
2338 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2339 .constrainAllUses(
TII,
TRI, RBI);
2346 return BuildCOPY(ResVReg, SrcPtr,
I);
2348 if ((SrcSC == SPIRV::StorageClass::Function &&
2349 DstSC == SPIRV::StorageClass::Private) ||
2350 (DstSC == SPIRV::StorageClass::Function &&
2351 SrcSC == SPIRV::StorageClass::Private))
2352 return BuildCOPY(ResVReg, SrcPtr,
I);
2356 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2359 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2362 SPIRVTypeInst GenericPtrTy =
2381 return selectUnOp(ResVReg, ResType,
I,
2382 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2384 return selectUnOp(ResVReg, ResType,
I,
2385 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2387 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2389 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2399 return SPIRV::OpFOrdEqual;
2401 return SPIRV::OpFOrdGreaterThanEqual;
2403 return SPIRV::OpFOrdGreaterThan;
2405 return SPIRV::OpFOrdLessThanEqual;
2407 return SPIRV::OpFOrdLessThan;
2409 return SPIRV::OpFOrdNotEqual;
2411 return SPIRV::OpOrdered;
2413 return SPIRV::OpFUnordEqual;
2415 return SPIRV::OpFUnordGreaterThanEqual;
2417 return SPIRV::OpFUnordGreaterThan;
2419 return SPIRV::OpFUnordLessThanEqual;
2421 return SPIRV::OpFUnordLessThan;
2423 return SPIRV::OpFUnordNotEqual;
2425 return SPIRV::OpUnordered;
2435 return SPIRV::OpIEqual;
2437 return SPIRV::OpINotEqual;
2439 return SPIRV::OpSGreaterThanEqual;
2441 return SPIRV::OpSGreaterThan;
2443 return SPIRV::OpSLessThanEqual;
2445 return SPIRV::OpSLessThan;
2447 return SPIRV::OpUGreaterThanEqual;
2449 return SPIRV::OpUGreaterThan;
2451 return SPIRV::OpULessThanEqual;
2453 return SPIRV::OpULessThan;
2462 return SPIRV::OpPtrEqual;
2464 return SPIRV::OpPtrNotEqual;
2475 return SPIRV::OpLogicalEqual;
2477 return SPIRV::OpLogicalNotEqual;
2511bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2512 SPIRVTypeInst ResType,
2514 unsigned OpAnyOrAll)
const {
2515 assert(
I.getNumOperands() == 3);
2516 assert(
I.getOperand(2).isReg());
2518 Register InputRegister =
I.getOperand(2).getReg();
2521 assert(InputType &&
"VReg has no type assigned");
2524 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2525 if (IsBoolTy && !IsVectorTy) {
2526 assert(ResVReg ==
I.getOperand(0).getReg());
2527 return BuildCOPY(ResVReg, InputRegister,
I);
2531 unsigned SpirvNotEqualId =
2532 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2534 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2539 IsBoolTy ? InputRegister
2547 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2549 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2566bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2567 SPIRVTypeInst ResType,
2568 MachineInstr &
I)
const {
2569 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2572bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2573 SPIRVTypeInst ResType,
2574 MachineInstr &
I)
const {
2575 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2579bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2580 SPIRVTypeInst ResType,
2581 MachineInstr &
I)
const {
2582 assert(
I.getNumOperands() == 4);
2583 assert(
I.getOperand(2).isReg());
2584 assert(
I.getOperand(3).isReg());
2586 [[maybe_unused]] SPIRVTypeInst VecType =
2591 "dot product requires a vector of at least 2 components");
2593 [[maybe_unused]] SPIRVTypeInst EltType =
2602 .
addUse(
I.getOperand(2).getReg())
2603 .
addUse(
I.getOperand(3).getReg())
2608bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2609 SPIRVTypeInst ResType,
2612 assert(
I.getNumOperands() == 4);
2613 assert(
I.getOperand(2).isReg());
2614 assert(
I.getOperand(3).isReg());
2617 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2621 .
addUse(
I.getOperand(2).getReg())
2622 .
addUse(
I.getOperand(3).getReg())
2629bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2630 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2631 assert(
I.getNumOperands() == 4);
2632 assert(
I.getOperand(2).isReg());
2633 assert(
I.getOperand(3).isReg());
2637 Register Vec0 =
I.getOperand(2).getReg();
2638 Register Vec1 =
I.getOperand(3).getReg();
2642 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2651 "dot product requires a vector of at least 2 components");
2654 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2664 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2675 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2687bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2688 SPIRVTypeInst ResType,
2689 MachineInstr &
I)
const {
2691 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2694 .
addUse(
I.getOperand(2).getReg())
2699bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2700 SPIRVTypeInst ResType,
2701 MachineInstr &
I)
const {
2703 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2706 .
addUse(
I.getOperand(2).getReg())
2711template <
bool Signed>
2712bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2713 SPIRVTypeInst ResType,
2714 MachineInstr &
I)
const {
2715 assert(
I.getNumOperands() == 5);
2716 assert(
I.getOperand(2).isReg());
2717 assert(
I.getOperand(3).isReg());
2718 assert(
I.getOperand(4).isReg());
2721 Register Acc =
I.getOperand(2).getReg();
2725 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2727 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
2732 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
2735 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2747template <
bool Signed>
2748bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
2749 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2750 assert(
I.getNumOperands() == 5);
2751 assert(
I.getOperand(2).isReg());
2752 assert(
I.getOperand(3).isReg());
2753 assert(
I.getOperand(4).isReg());
2756 Register Acc =
I.getOperand(2).getReg();
2762 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
2766 for (
unsigned i = 0; i < 4; i++) {
2789 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
2809 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2824bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
2825 SPIRVTypeInst ResType,
2826 MachineInstr &
I)
const {
2827 assert(
I.getNumOperands() == 3);
2828 assert(
I.getOperand(2).isReg());
2830 Register VZero = buildZerosValF(ResType,
I);
2831 Register VOne = buildOnesValF(ResType,
I);
2833 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
2836 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2838 .
addUse(
I.getOperand(2).getReg())
2845bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
2846 SPIRVTypeInst ResType,
2847 MachineInstr &
I)
const {
2848 assert(
I.getNumOperands() == 3);
2849 assert(
I.getOperand(2).isReg());
2851 Register InputRegister =
I.getOperand(2).getReg();
2853 auto &
DL =
I.getDebugLoc();
2863 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
2865 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
2873 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2878 if (NeedsConversion) {
2879 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
2890bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
2891 SPIRVTypeInst ResType,
2893 unsigned Opcode)
const {
2897 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2903 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
2904 BMI.addUse(
I.getOperand(J).getReg());
2911bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
2914 bool WithGroupSync)
const {
2916 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
2918 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
2920 assert(((Scope != SPIRV::Scope::Workgroup) ||
2921 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
2922 "Workgroup Scope must set WorkGroupMemory semantic "
2923 "in Barrier instruction");
2925 assert(((Scope != SPIRV::Scope::Device) ||
2926 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
2927 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
2928 "Device Scope must set UniformMemory and ImageMemory semantic "
2929 "in Barrier instruction");
2931 Register MemSemReg = buildI32Constant(MemSem,
I);
2932 Register ScopeReg = buildI32Constant(Scope,
I);
2938 if (WithGroupSync) {
2939 MI.addUse(ScopeReg);
2942 MI.addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
2946bool SPIRVInstructionSelector::selectWaveActiveCountBits(
2947 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2952 if (!selectWaveOpInst(BallotReg, BallotType,
I,
2953 SPIRV::OpGroupNonUniformBallot))
2958 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2963 .
addImm(SPIRV::GroupOperation::Reduce)
2972 if (
Type->getOpcode() != SPIRV::OpTypeVector)
2976 return Type->getOperand(2).getImm();
2979bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
2980 SPIRVTypeInst ResType,
2981 MachineInstr &
I)
const {
2986 Register InputReg =
I.getOperand(2).getReg();
2991 bool IsVector = NumElems > 1;
2994 SPIRVTypeInst ElemInputType = InputType;
2995 SPIRVTypeInst ElemBoolType = ResType;
3008 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3009 SPIRV::OpGroupNonUniformAllEqual);
3014 ElementResults.
reserve(NumElems);
3016 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3029 ElemInput = Extracted;
3035 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3046 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3057bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3058 SPIRVTypeInst ResType,
3059 MachineInstr &
I)
const {
3061 assert(
I.getNumOperands() == 3);
3063 auto Op =
I.getOperand(2);
3075 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3097 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3101 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3108bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3109 SPIRVTypeInst ResType,
3111 bool IsUnsigned)
const {
3112 return selectWaveReduce(
3113 ResVReg, ResType,
I, IsUnsigned,
3114 [&](
Register InputRegister,
bool IsUnsigned) {
3115 const bool IsFloatTy =
3117 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3118 : SPIRV::OpGroupNonUniformSMax;
3119 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3123bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3124 SPIRVTypeInst ResType,
3126 bool IsUnsigned)
const {
3127 return selectWaveReduce(
3128 ResVReg, ResType,
I, IsUnsigned,
3129 [&](
Register InputRegister,
bool IsUnsigned) {
3130 const bool IsFloatTy =
3132 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3133 : SPIRV::OpGroupNonUniformSMin;
3134 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3138bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3139 SPIRVTypeInst ResType,
3140 MachineInstr &
I)
const {
3141 return selectWaveReduce(ResVReg, ResType,
I,
false,
3142 [&](
Register InputRegister,
bool IsUnsigned) {
3144 InputRegister, SPIRV::OpTypeFloat);
3145 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3146 : SPIRV::OpGroupNonUniformIAdd;
3150bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3151 SPIRVTypeInst ResType,
3152 MachineInstr &
I)
const {
3153 return selectWaveReduce(ResVReg, ResType,
I,
false,
3154 [&](
Register InputRegister,
bool IsUnsigned) {
3156 InputRegister, SPIRV::OpTypeFloat);
3157 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3158 : SPIRV::OpGroupNonUniformIMul;
3162template <
typename PickOpcodeFn>
3163bool SPIRVInstructionSelector::selectWaveReduce(
3164 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3165 PickOpcodeFn &&PickOpcode)
const {
3166 assert(
I.getNumOperands() == 3);
3167 assert(
I.getOperand(2).isReg());
3169 Register InputRegister =
I.getOperand(2).getReg();
3176 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3182 .
addImm(SPIRV::GroupOperation::Reduce)
3183 .
addUse(
I.getOperand(2).getReg())
3188bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3189 SPIRVTypeInst ResType,
3191 unsigned Opcode)
const {
3192 return selectWaveReduce(
3193 ResVReg, ResType,
I,
false,
3194 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3197bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3198 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3199 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3200 [&](
Register InputRegister,
bool IsUnsigned) {
3202 InputRegister, SPIRV::OpTypeFloat);
3204 ? SPIRV::OpGroupNonUniformFAdd
3205 : SPIRV::OpGroupNonUniformIAdd;
3209bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3210 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3211 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3212 [&](
Register InputRegister,
bool IsUnsigned) {
3214 InputRegister, SPIRV::OpTypeFloat);
3216 ? SPIRV::OpGroupNonUniformFMul
3217 : SPIRV::OpGroupNonUniformIMul;
3221template <
typename PickOpcodeFn>
3222bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3223 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3224 PickOpcodeFn &&PickOpcode)
const {
3225 assert(
I.getNumOperands() == 3);
3226 assert(
I.getOperand(2).isReg());
3228 Register InputRegister =
I.getOperand(2).getReg();
3235 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3241 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3242 .
addUse(
I.getOperand(2).getReg())
3247bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3248 SPIRVTypeInst ResType,
3251 assert(
I.getNumOperands() == 3);
3252 assert(
I.getOperand(2).isReg());
3254 Register InputRegister =
I.getOperand(2).getReg();
3260 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3271bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3272 SPIRVTypeInst ResType,
3277 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3282 : SPIRV::OpUConvert;
3286 ShiftOp = SPIRV::OpShiftRightLogicalV;
3291 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3292 TII.get(SPIRV::OpConstantComposite))
3295 for (
unsigned It = 0; It <
N; ++It)
3299 ShiftConst = CompositeReg;
3304 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3309 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3314 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3319 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3322bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
3323 SPIRVTypeInst ResType,
3327 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3335bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3336 SPIRVTypeInst ResType,
3337 MachineInstr &
I)
const {
3338 Register OpReg =
I.getOperand(1).getReg();
3345 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3347 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3352 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3363 unsigned AndOp = SPIRV::OpBitwiseAndS;
3364 unsigned OrOp = SPIRV::OpBitwiseOrS;
3365 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
3366 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
3368 AndOp = SPIRV::OpBitwiseAndV;
3369 OrOp = SPIRV::OpBitwiseOrV;
3370 ShlOp = SPIRV::OpShiftLeftLogicalV;
3371 ShrOp = SPIRV::OpShiftRightLogicalV;
3377 const unsigned Shift) ->
Register {
3385 Register MaskReg = CreateConst(Mask);
3386 Register ShiftReg = CreateConst(Shift);
3393 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
3394 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
3395 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
3396 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
3397 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
3405 uint64_t
Mask = ~0ull;
3406 while ((Shift >>= 1) > 0) {
3413 return BuildCOPY(ResVReg, Result,
I);
3416bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3417 SPIRVTypeInst ResType,
3418 MachineInstr &
I)
const {
3424 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3426 Register OpReg =
I.getOperand(1).getReg();
3427 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
3428 if (
Def->getOpcode() == TargetOpcode::COPY)
3431 switch (
Def->getOpcode()) {
3432 case SPIRV::ASSIGN_TYPE:
3433 if (MachineInstr *AssignToDef =
3435 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3436 Reg =
Def->getOperand(2).getReg();
3439 case SPIRV::OpUndef:
3440 Reg =
Def->getOperand(1).getReg();
3443 unsigned DestOpCode;
3445 DestOpCode = SPIRV::OpConstantNull;
3447 DestOpCode = TargetOpcode::COPY;
3450 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3451 .
addDef(
I.getOperand(0).getReg())
3459bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3460 SPIRVTypeInst ResType,
3461 MachineInstr &
I)
const {
3463 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3465 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3469 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3474 for (
unsigned i =
I.getNumExplicitDefs();
3475 i <
I.getNumExplicitOperands() && IsConst; ++i)
3479 if (!IsConst &&
N < 2)
3481 "There must be at least two constituent operands in a vector");
3484 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3485 TII.get(IsConst ? SPIRV::OpConstantComposite
3486 : SPIRV::OpCompositeConstruct))
3489 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3490 MIB.
addUse(
I.getOperand(i).getReg());
3495bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3496 SPIRVTypeInst ResType,
3497 MachineInstr &
I)
const {
3499 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3501 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3507 if (!
I.getOperand(
OpIdx).isReg())
3514 if (!IsConst &&
N < 2)
3516 "There must be at least two constituent operands in a vector");
3519 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3520 TII.get(IsConst ? SPIRV::OpConstantComposite
3521 : SPIRV::OpCompositeConstruct))
3524 for (
unsigned i = 0; i <
N; ++i)
3530bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3531 SPIRVTypeInst ResType,
3532 MachineInstr &
I)
const {
3537 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3539 Opcode = SPIRV::OpDemoteToHelperInvocation;
3541 Opcode = SPIRV::OpKill;
3543 if (MachineInstr *NextI =
I.getNextNode()) {
3545 NextI->eraseFromParent();
3555bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3556 SPIRVTypeInst ResType,
unsigned CmpOpc,
3557 MachineInstr &
I)
const {
3558 Register Cmp0 =
I.getOperand(2).getReg();
3559 Register Cmp1 =
I.getOperand(3).getReg();
3562 "CMP operands should have the same type");
3563 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3573bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3574 SPIRVTypeInst ResType,
3575 MachineInstr &
I)
const {
3576 auto Pred =
I.getOperand(1).getPredicate();
3579 Register CmpOperand =
I.getOperand(2).getReg();
3586 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3590SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3591 SPIRVTypeInst ResType)
const {
3593 SPIRVTypeInst SpvI32Ty =
3596 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3603 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3606 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3609 .
addImm(APInt(32, Val).getZExtValue());
3611 GR.
add(ConstInt,
MI);
3616bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
3617 SPIRVTypeInst ResType,
3618 MachineInstr &
I)
const {
3620 return selectCmp(ResVReg, ResType, CmpOp,
I);
3623bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
3624 SPIRVTypeInst ResType,
3625 MachineInstr &
I)
const {
3627 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
3634 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
3635 ResType->
getOpcode() != SPIRV::OpTypeFloat)
3638 MachineIRBuilder MIRBuilder(
I);
3640 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
3646 "only float operands supported by GLSL extended math");
3649 MIRBuilder, SpirvScalarType);
3651 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
3652 ? SPIRV::OpVectorTimesScalar
3655 if (!selectOpWithSrcs(ArgReg, ResType,
I,
3656 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
3658 if (!selectExtInst(ResVReg, ResType,
I,
3659 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
3669Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
3670 MachineInstr &
I)
const {
3673 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3678bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
3684 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
3692 if (
Def->getOpcode() == SPIRV::OpConstantNull)
3695 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
3696 Def->getOpcode() == SPIRV::OpConstantI)
3709 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
3710 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
3712 Intrinsic::spv_const_composite)) {
3713 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
3714 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
3715 if (!IsZero(
Def->getOperand(i).getReg()))
3724Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
3725 MachineInstr &
I)
const {
3729 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3734Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
3735 MachineInstr &
I)
const {
3739 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3745 SPIRVTypeInst ResType,
3746 MachineInstr &
I)
const {
3750 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3755bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
3756 SPIRVTypeInst ResType,
3757 MachineInstr &
I)
const {
3758 Register SelectFirstArg =
I.getOperand(2).getReg();
3759 Register SelectSecondArg =
I.getOperand(3).getReg();
3768 SPIRV::OpTypeVector;
3775 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
3776 }
else if (IsPtrTy) {
3777 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
3779 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
3783 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
3784 }
else if (IsPtrTy) {
3785 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
3787 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3790 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3793 .
addUse(
I.getOperand(1).getReg())
3802bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
3803 SPIRVTypeInst ResType,
3805 MachineInstr &InsertAt,
3806 bool IsSigned)
const {
3808 Register ZeroReg = buildZerosVal(ResType, InsertAt);
3809 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
3810 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
3812 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3824bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
3825 SPIRVTypeInst ResType,
3826 MachineInstr &
I,
bool IsSigned,
3827 unsigned Opcode)
const {
3828 Register SrcReg =
I.getOperand(1).getReg();
3834 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
3839 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
3841 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
3844bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
3845 SPIRVTypeInst ResType, MachineInstr &
I,
3846 bool IsSigned)
const {
3847 Register SrcReg =
I.getOperand(1).getReg();
3849 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
3853 if (ResType == SrcType)
3854 return BuildCOPY(ResVReg, SrcReg,
I);
3856 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3857 return selectUnOp(ResVReg, ResType,
I, Opcode);
3860bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
3861 SPIRVTypeInst ResType,
3863 bool IsSigned)
const {
3864 MachineIRBuilder MIRBuilder(
I);
3865 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
3880 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
3883 .
addUse(
I.getOperand(1).getReg())
3884 .
addUse(
I.getOperand(2).getReg())
3890 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
3893 .
addUse(
I.getOperand(1).getReg())
3894 .
addUse(
I.getOperand(2).getReg())
3902 unsigned SelectOpcode =
3903 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
3908 .
addUse(buildOnesVal(
true, ResType,
I))
3909 .
addUse(buildZerosVal(ResType,
I))
3916 .
addUse(buildOnesVal(
false, ResType,
I))
3921bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
3924 SPIRVTypeInst IntTy,
3925 SPIRVTypeInst BoolTy)
const {
3928 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
3929 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
3931 Register One = buildOnesVal(
false, IntTy,
I);
3939 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
3948bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
3949 SPIRVTypeInst ResType,
3950 MachineInstr &
I)
const {
3951 Register IntReg =
I.getOperand(1).getReg();
3954 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
3955 if (ArgType == ResType)
3956 return BuildCOPY(ResVReg, IntReg,
I);
3958 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3959 return selectUnOp(ResVReg, ResType,
I, Opcode);
3962bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
3963 SPIRVTypeInst ResType,
3964 MachineInstr &
I)
const {
3965 unsigned Opcode =
I.getOpcode();
3966 unsigned TpOpcode = ResType->
getOpcode();
3968 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
3969 assert(Opcode == TargetOpcode::G_CONSTANT &&
3970 I.getOperand(1).getCImm()->isZero());
3971 MachineBasicBlock &DepMBB =
I.getMF()->front();
3974 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
3981 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
3984bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
3985 SPIRVTypeInst ResType,
3986 MachineInstr &
I)
const {
3987 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3994bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
3995 SPIRVTypeInst ResType,
3996 MachineInstr &
I)
const {
3998 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4002 .
addUse(
I.getOperand(3).getReg())
4004 .
addUse(
I.getOperand(2).getReg());
4005 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4011bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4012 SPIRVTypeInst ResType,
4013 MachineInstr &
I)
const {
4014 Type *MaybeResTy =
nullptr;
4019 "Expected aggregate type for extractv instruction");
4021 SPIRV::AccessQualifier::ReadWrite,
false);
4025 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4028 .
addUse(
I.getOperand(2).getReg());
4029 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4035bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4036 SPIRVTypeInst ResType,
4037 MachineInstr &
I)
const {
4038 if (
getImm(
I.getOperand(4), MRI))
4039 return selectInsertVal(ResVReg, ResType,
I);
4041 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4044 .
addUse(
I.getOperand(2).getReg())
4045 .
addUse(
I.getOperand(3).getReg())
4046 .
addUse(
I.getOperand(4).getReg())
4051bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4052 SPIRVTypeInst ResType,
4053 MachineInstr &
I)
const {
4054 if (
getImm(
I.getOperand(3), MRI))
4055 return selectExtractVal(ResVReg, ResType,
I);
4057 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4060 .
addUse(
I.getOperand(2).getReg())
4061 .
addUse(
I.getOperand(3).getReg())
4066bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4067 SPIRVTypeInst ResType,
4068 MachineInstr &
I)
const {
4069 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4075 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4076 : SPIRV::OpAccessChain)
4077 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4078 :
SPIRV::OpPtrAccessChain);
4080 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4084 .
addUse(
I.getOperand(3).getReg());
4086 (Opcode == SPIRV::OpPtrAccessChain ||
4087 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4088 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4089 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4092 const unsigned StartingIndex =
4093 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4096 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4097 Res.addUse(
I.getOperand(i).getReg());
4098 Res.constrainAllUses(
TII,
TRI, RBI);
4103bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4105 unsigned Lim =
I.getNumExplicitOperands();
4106 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4107 Register OpReg =
I.getOperand(i).getReg();
4108 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4110 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4111 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4112 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4119 MachineFunction *MF =
I.getMF();
4131 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4132 TII.get(SPIRV::OpSpecConstantOp))
4135 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4137 GR.
add(OpDefine, MIB);
4143bool SPIRVInstructionSelector::selectDerivativeInst(
4144 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4145 const unsigned DPdOpCode)
const {
4148 errorIfInstrOutsideShader(
I);
4153 Register SrcReg =
I.getOperand(2).getReg();
4158 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4161 .
addUse(
I.getOperand(2).getReg());
4163 MachineIRBuilder MIRBuilder(
I);
4166 if (componentCount != 1)
4170 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4174 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4179 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4184 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4192bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4193 SPIRVTypeInst ResType,
4194 MachineInstr &
I)
const {
4198 case Intrinsic::spv_load:
4199 return selectLoad(ResVReg, ResType,
I);
4200 case Intrinsic::spv_store:
4201 return selectStore(
I);
4202 case Intrinsic::spv_extractv:
4203 return selectExtractVal(ResVReg, ResType,
I);
4204 case Intrinsic::spv_insertv:
4205 return selectInsertVal(ResVReg, ResType,
I);
4206 case Intrinsic::spv_extractelt:
4207 return selectExtractElt(ResVReg, ResType,
I);
4208 case Intrinsic::spv_insertelt:
4209 return selectInsertElt(ResVReg, ResType,
I);
4210 case Intrinsic::spv_gep:
4211 return selectGEP(ResVReg, ResType,
I);
4212 case Intrinsic::spv_bitcast: {
4213 Register OpReg =
I.getOperand(2).getReg();
4214 SPIRVTypeInst OpType =
4218 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4220 case Intrinsic::spv_unref_global:
4221 case Intrinsic::spv_init_global: {
4222 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4227 Register GVarVReg =
MI->getOperand(0).getReg();
4228 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4233 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4235 MI->eraseFromParent();
4239 case Intrinsic::spv_undef: {
4240 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4246 case Intrinsic::spv_named_boolean_spec_constant: {
4247 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4248 : SPIRV::OpSpecConstantFalse;
4250 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4251 .
addDef(
I.getOperand(0).getReg())
4254 unsigned SpecId =
I.getOperand(2).getImm();
4256 SPIRV::Decoration::SpecId, {SpecId});
4260 case Intrinsic::spv_const_composite: {
4262 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4268 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4270 std::function<bool(
Register)> HasSpecConstOperand =
4280 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
4281 J < Def->getNumExplicitOperands(); ++J) {
4282 if (
Def->getOperand(J).isReg() &&
4283 HasSpecConstOperand(
Def->getOperand(J).getReg()))
4289 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
4290 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
4291 : SPIRV::OpConstantComposite;
4292 unsigned ContinuedOpc = HasSpecConst
4293 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
4294 : SPIRV::OpConstantCompositeContinuedINTEL;
4295 MachineIRBuilder MIR(
I);
4297 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
4299 for (
auto *Instr : Instructions) {
4300 Instr->setDebugLoc(
I.getDebugLoc());
4305 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4312 case Intrinsic::spv_assign_name: {
4313 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4314 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4315 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4316 i <
I.getNumExplicitOperands(); ++i) {
4317 MIB.
addImm(
I.getOperand(i).getImm());
4322 case Intrinsic::spv_switch: {
4323 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4324 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4325 if (
I.getOperand(i).isReg())
4326 MIB.
addReg(
I.getOperand(i).getReg());
4327 else if (
I.getOperand(i).isCImm())
4328 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4329 else if (
I.getOperand(i).isMBB())
4330 MIB.
addMBB(
I.getOperand(i).getMBB());
4337 case Intrinsic::spv_loop_merge: {
4338 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4339 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4340 if (
I.getOperand(i).isMBB())
4341 MIB.
addMBB(
I.getOperand(i).getMBB());
4348 case Intrinsic::spv_loop_control_intel: {
4350 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4351 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4356 case Intrinsic::spv_selection_merge: {
4358 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4359 assert(
I.getOperand(1).isMBB() &&
4360 "operand 1 to spv_selection_merge must be a basic block");
4361 MIB.
addMBB(
I.getOperand(1).getMBB());
4362 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4366 case Intrinsic::spv_cmpxchg:
4367 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4368 case Intrinsic::spv_unreachable:
4369 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4372 case Intrinsic::spv_alloca:
4373 return selectFrameIndex(ResVReg, ResType,
I);
4374 case Intrinsic::spv_alloca_array:
4375 return selectAllocaArray(ResVReg, ResType,
I);
4376 case Intrinsic::spv_assume:
4378 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4379 .
addUse(
I.getOperand(1).getReg())
4384 case Intrinsic::spv_expect:
4386 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4389 .
addUse(
I.getOperand(2).getReg())
4390 .
addUse(
I.getOperand(3).getReg())
4395 case Intrinsic::arithmetic_fence:
4396 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4397 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4400 .
addUse(
I.getOperand(2).getReg())
4404 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4406 case Intrinsic::spv_thread_id:
4412 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4414 case Intrinsic::spv_thread_id_in_group:
4420 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4422 case Intrinsic::spv_group_id:
4428 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4430 case Intrinsic::spv_flattened_thread_id_in_group:
4437 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4439 case Intrinsic::spv_workgroup_size:
4440 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4442 case Intrinsic::spv_global_size:
4443 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4445 case Intrinsic::spv_global_offset:
4446 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4448 case Intrinsic::spv_num_workgroups:
4449 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4451 case Intrinsic::spv_subgroup_size:
4452 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4454 case Intrinsic::spv_num_subgroups:
4455 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4457 case Intrinsic::spv_subgroup_id:
4458 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4459 case Intrinsic::spv_subgroup_local_invocation_id:
4460 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4461 ResVReg, ResType,
I);
4462 case Intrinsic::spv_subgroup_max_size:
4463 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4465 case Intrinsic::spv_fdot:
4466 return selectFloatDot(ResVReg, ResType,
I);
4467 case Intrinsic::spv_udot:
4468 case Intrinsic::spv_sdot:
4469 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4471 return selectIntegerDot(ResVReg, ResType,
I,
4472 IID == Intrinsic::spv_sdot);
4473 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4474 case Intrinsic::spv_dot4add_i8packed:
4475 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4477 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4478 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4479 case Intrinsic::spv_dot4add_u8packed:
4480 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4482 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4483 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4484 case Intrinsic::spv_all:
4485 return selectAll(ResVReg, ResType,
I);
4486 case Intrinsic::spv_any:
4487 return selectAny(ResVReg, ResType,
I);
4488 case Intrinsic::spv_cross:
4489 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4490 case Intrinsic::spv_distance:
4491 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4492 case Intrinsic::spv_lerp:
4493 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4494 case Intrinsic::spv_length:
4495 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4496 case Intrinsic::spv_degrees:
4497 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4498 case Intrinsic::spv_faceforward:
4499 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4500 case Intrinsic::spv_frac:
4501 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4502 case Intrinsic::spv_isinf:
4503 return selectOpIsInf(ResVReg, ResType,
I);
4504 case Intrinsic::spv_isnan:
4505 return selectOpIsNan(ResVReg, ResType,
I);
4506 case Intrinsic::spv_normalize:
4507 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4508 case Intrinsic::spv_refract:
4509 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4510 case Intrinsic::spv_reflect:
4511 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4512 case Intrinsic::spv_rsqrt:
4513 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4514 case Intrinsic::spv_sign:
4515 return selectSign(ResVReg, ResType,
I);
4516 case Intrinsic::spv_smoothstep:
4517 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4518 case Intrinsic::spv_firstbituhigh:
4519 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4520 case Intrinsic::spv_firstbitshigh:
4521 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4522 case Intrinsic::spv_firstbitlow:
4523 return selectFirstBitLow(ResVReg, ResType,
I);
4524 case Intrinsic::spv_group_memory_barrier:
4525 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4526 SPIRV::MemorySemantics::WorkgroupMemory,
4528 case Intrinsic::spv_group_memory_barrier_with_group_sync:
4529 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4530 SPIRV::MemorySemantics::WorkgroupMemory,
4532 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4533 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4534 SPIRV::StorageClass::StorageClass ResSC =
4538 "Generic storage class");
4539 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4547 case Intrinsic::spv_lifetime_start:
4548 case Intrinsic::spv_lifetime_end: {
4549 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
4550 : SPIRV::OpLifetimeStop;
4551 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
4552 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
4561 case Intrinsic::spv_saturate:
4562 return selectSaturate(ResVReg, ResType,
I);
4563 case Intrinsic::spv_nclamp:
4564 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
4565 case Intrinsic::spv_uclamp:
4566 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
4567 case Intrinsic::spv_sclamp:
4568 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
4569 case Intrinsic::spv_subgroup_prefix_bit_count:
4570 return selectWavePrefixBitCount(ResVReg, ResType,
I);
4571 case Intrinsic::spv_wave_active_countbits:
4572 return selectWaveActiveCountBits(ResVReg, ResType,
I);
4573 case Intrinsic::spv_wave_all_equal:
4574 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
4575 case Intrinsic::spv_wave_all:
4576 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
4577 case Intrinsic::spv_wave_any:
4578 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
4579 case Intrinsic::spv_subgroup_ballot:
4580 return selectWaveOpInst(ResVReg, ResType,
I,
4581 SPIRV::OpGroupNonUniformBallot);
4582 case Intrinsic::spv_wave_is_first_lane:
4583 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
4584 case Intrinsic::spv_wave_reduce_or:
4585 return selectWaveReduceOp(ResVReg, ResType,
I,
4586 SPIRV::OpGroupNonUniformBitwiseOr);
4587 case Intrinsic::spv_wave_reduce_xor:
4588 return selectWaveReduceOp(ResVReg, ResType,
I,
4589 SPIRV::OpGroupNonUniformBitwiseXor);
4590 case Intrinsic::spv_wave_reduce_and:
4591 return selectWaveReduceOp(ResVReg, ResType,
I,
4592 SPIRV::OpGroupNonUniformBitwiseAnd);
4593 case Intrinsic::spv_wave_reduce_umax:
4594 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
4595 case Intrinsic::spv_wave_reduce_max:
4596 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
4597 case Intrinsic::spv_wave_reduce_umin:
4598 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
4599 case Intrinsic::spv_wave_reduce_min:
4600 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
4601 case Intrinsic::spv_wave_reduce_sum:
4602 return selectWaveReduceSum(ResVReg, ResType,
I);
4603 case Intrinsic::spv_wave_product:
4604 return selectWaveReduceProduct(ResVReg, ResType,
I);
4605 case Intrinsic::spv_wave_readlane:
4606 return selectWaveOpInst(ResVReg, ResType,
I,
4607 SPIRV::OpGroupNonUniformShuffle);
4608 case Intrinsic::spv_wave_prefix_sum:
4609 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
4610 case Intrinsic::spv_wave_prefix_product:
4611 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
4612 case Intrinsic::spv_quad_read_across_x: {
4613 return selectQuadSwap(ResVReg, ResType,
I, 0);
4615 case Intrinsic::spv_quad_read_across_y: {
4616 return selectQuadSwap(ResVReg, ResType,
I, 1);
4618 case Intrinsic::spv_step:
4619 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
4620 case Intrinsic::spv_radians:
4621 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
4625 case Intrinsic::instrprof_increment:
4626 case Intrinsic::instrprof_increment_step:
4627 case Intrinsic::instrprof_value_profile:
4630 case Intrinsic::spv_value_md:
4632 case Intrinsic::spv_resource_handlefrombinding: {
4633 return selectHandleFromBinding(ResVReg, ResType,
I);
4635 case Intrinsic::spv_resource_counterhandlefrombinding:
4636 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
4637 case Intrinsic::spv_resource_updatecounter:
4638 return selectUpdateCounter(ResVReg, ResType,
I);
4639 case Intrinsic::spv_resource_store_typedbuffer: {
4640 return selectImageWriteIntrinsic(
I);
4642 case Intrinsic::spv_resource_load_typedbuffer: {
4643 return selectReadImageIntrinsic(ResVReg, ResType,
I);
4645 case Intrinsic::spv_resource_load_level: {
4646 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
4648 case Intrinsic::spv_resource_getdimensions_x:
4649 case Intrinsic::spv_resource_getdimensions_xy:
4650 case Intrinsic::spv_resource_getdimensions_xyz: {
4651 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
4653 case Intrinsic::spv_resource_getdimensions_levels_x:
4654 case Intrinsic::spv_resource_getdimensions_levels_xy:
4655 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
4656 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
4658 case Intrinsic::spv_resource_getdimensions_ms_xy:
4659 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
4660 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
4662 case Intrinsic::spv_resource_calculate_lod:
4663 case Intrinsic::spv_resource_calculate_lod_unclamped:
4664 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
4665 case Intrinsic::spv_resource_sample:
4666 case Intrinsic::spv_resource_sample_clamp:
4667 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
4668 case Intrinsic::spv_resource_samplebias:
4669 case Intrinsic::spv_resource_samplebias_clamp:
4670 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
4671 case Intrinsic::spv_resource_samplegrad:
4672 case Intrinsic::spv_resource_samplegrad_clamp:
4673 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
4674 case Intrinsic::spv_resource_samplelevel:
4675 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
4676 case Intrinsic::spv_resource_samplecmp:
4677 case Intrinsic::spv_resource_samplecmp_clamp:
4678 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
4679 case Intrinsic::spv_resource_samplecmplevelzero:
4680 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
4681 case Intrinsic::spv_resource_gather:
4682 case Intrinsic::spv_resource_gather_cmp:
4683 return selectGatherIntrinsic(ResVReg, ResType,
I);
4684 case Intrinsic::spv_resource_getpointer: {
4685 return selectResourceGetPointer(ResVReg, ResType,
I);
4687 case Intrinsic::spv_pushconstant_getpointer: {
4688 return selectPushConstantGetPointer(ResVReg, ResType,
I);
4690 case Intrinsic::spv_discard: {
4691 return selectDiscard(ResVReg, ResType,
I);
4693 case Intrinsic::spv_resource_nonuniformindex: {
4694 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
4696 case Intrinsic::spv_unpackhalf2x16: {
4697 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
4699 case Intrinsic::spv_packhalf2x16: {
4700 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
4702 case Intrinsic::spv_ddx:
4703 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
4704 case Intrinsic::spv_ddy:
4705 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
4706 case Intrinsic::spv_ddx_coarse:
4707 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
4708 case Intrinsic::spv_ddy_coarse:
4709 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
4710 case Intrinsic::spv_ddx_fine:
4711 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
4712 case Intrinsic::spv_ddy_fine:
4713 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
4714 case Intrinsic::spv_fwidth:
4715 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
4716 case Intrinsic::spv_masked_gather:
4717 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4718 return selectMaskedGather(ResVReg, ResType,
I);
4719 return diagnoseUnsupported(
4720 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
4721 case Intrinsic::spv_masked_scatter:
4722 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4723 return selectMaskedScatter(
I);
4724 return diagnoseUnsupported(
4725 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
4727 std::string DiagMsg;
4728 raw_string_ostream OS(DiagMsg);
4730 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
4737bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
4738 SPIRVTypeInst ResType,
4739 MachineInstr &
I)
const {
4742 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
4749bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
4750 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4752 assert(Intr.getIntrinsicID() ==
4753 Intrinsic::spv_resource_counterhandlefrombinding);
4756 Register MainHandleReg = Intr.getOperand(2).getReg();
4758 assert(MainHandleDef->getIntrinsicID() ==
4759 Intrinsic::spv_resource_handlefrombinding);
4763 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
4764 Register IndexReg = MainHandleDef->getOperand(5).getReg();
4765 std::string CounterName =
4770 MachineIRBuilder MIRBuilder(
I);
4772 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
4774 ArraySize, IndexReg, CounterName, MIRBuilder);
4776 return BuildCOPY(ResVReg, CounterVarReg,
I);
4779bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
4780 SPIRVTypeInst ResType,
4781 MachineInstr &
I)
const {
4783 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
4785 Register CounterHandleReg = Intr.getOperand(2).getReg();
4786 Register IncrReg = Intr.getOperand(3).getReg();
4793 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
4794 assert(CounterVarPointeeType &&
4795 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
4796 "Counter variable must be a struct");
4798 SPIRV::StorageClass::StorageBuffer &&
4799 "Counter variable must be in the storage buffer storage class");
4801 "Counter variable must have exactly 1 member in the struct");
4802 const SPIRVTypeInst MemberType =
4805 "Counter variable struct must have a single i32 member");
4809 MachineIRBuilder MIRBuilder(
I);
4811 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
4814 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
4820 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4823 .
addUse(CounterHandleReg)
4830 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
4833 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
4836 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
4845 return BuildCOPY(ResVReg, AtomicRes,
I);
4853 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
4861bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
4862 SPIRVTypeInst ResType,
4863 MachineInstr &
I)
const {
4871 Register ImageReg =
I.getOperand(2).getReg();
4879 Register IdxReg =
I.getOperand(3).getReg();
4881 MachineInstr &Pos =
I;
4883 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
4887bool SPIRVInstructionSelector::generateSampleImage(
4890 DebugLoc Loc, MachineInstr &Pos)
const {
4901 if (!loadHandleBeforePosition(NewSamplerReg,
4907 MachineIRBuilder MIRBuilder(Pos);
4920 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
4921 ImOps.Lod.has_value();
4922 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
4923 : SPIRV::OpImageSampleImplicitLod;
4925 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
4926 : SPIRV::OpImageSampleDrefImplicitLod;
4935 MIB.
addUse(*ImOps.Compare);
4937 uint32_t ImageOperands = 0;
4939 ImageOperands |= SPIRV::ImageOperand::Bias;
4941 ImageOperands |= SPIRV::ImageOperand::Lod;
4942 if (ImOps.GradX && ImOps.GradY)
4943 ImageOperands |= SPIRV::ImageOperand::Grad;
4944 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
4946 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4949 "Non-constant offsets are not supported in sample instructions.");
4953 ImageOperands |= SPIRV::ImageOperand::MinLod;
4955 if (ImageOperands != 0) {
4956 MIB.
addImm(ImageOperands);
4957 if (ImageOperands & SPIRV::ImageOperand::Bias)
4959 if (ImageOperands & SPIRV::ImageOperand::Lod)
4961 if (ImageOperands & SPIRV::ImageOperand::Grad) {
4962 MIB.
addUse(*ImOps.GradX);
4963 MIB.
addUse(*ImOps.GradY);
4966 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4967 MIB.
addUse(*ImOps.Offset);
4968 if (ImageOperands & SPIRV::ImageOperand::MinLod)
4969 MIB.
addUse(*ImOps.MinLod);
4976bool SPIRVInstructionSelector::selectImageQuerySize(
4978 std::optional<Register> LodReg)
const {
4980 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
4983 "ImageReg is not an image type.");
4985 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
4987 unsigned NumComponents = 0;
4989 case SPIRV::Dim::DIM_1D:
4990 case SPIRV::Dim::DIM_Buffer:
4991 NumComponents =
IsArray ? 2 : 1;
4993 case SPIRV::Dim::DIM_2D:
4994 case SPIRV::Dim::DIM_Cube:
4995 case SPIRV::Dim::DIM_Rect:
4996 NumComponents =
IsArray ? 3 : 2;
4998 case SPIRV::Dim::DIM_3D:
5002 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5007 SPIRVTypeInst ResType =
5012 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5022bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5023 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5024 Register ImageReg =
I.getOperand(2).getReg();
5031 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5034bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5035 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5036 Register ImageReg =
I.getOperand(2).getReg();
5045 Register LodReg =
I.getOperand(3).getReg();
5048 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5050 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5057 TII.get(SPIRV::OpImageQueryLevels))
5064 TII.get(SPIRV::OpCompositeConstruct))
5074bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5075 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5076 Register ImageReg =
I.getOperand(2).getReg();
5087 "OpImageQuerySamples requires a multisampled image");
5089 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5097 TII.get(SPIRV::OpImageQuerySamples))
5104 TII.get(SPIRV::OpCompositeConstruct))
5114bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5115 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5116 Register ImageReg =
I.getOperand(2).getReg();
5117 Register SamplerReg =
I.getOperand(3).getReg();
5118 Register CoordinateReg =
I.getOperand(4).getReg();
5134 if (!loadHandleBeforePosition(
5139 MachineIRBuilder MIRBuilder(
I);
5145 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5155 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
5162 unsigned ExtractedIndex =
5164 Intrinsic::spv_resource_calculate_lod_unclamped
5168 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5169 TII.get(SPIRV::OpCompositeExtract))
5179bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
5180 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5181 Register ImageReg =
I.getOperand(2).getReg();
5182 Register SamplerReg =
I.getOperand(3).getReg();
5183 Register CoordinateReg =
I.getOperand(4).getReg();
5184 ImageOperands ImOps;
5185 if (
I.getNumOperands() > 5)
5186 ImOps.Offset =
I.getOperand(5).getReg();
5187 if (
I.getNumOperands() > 6)
5188 ImOps.MinLod =
I.getOperand(6).getReg();
5189 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5190 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5193bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5194 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5195 Register ImageReg =
I.getOperand(2).getReg();
5196 Register SamplerReg =
I.getOperand(3).getReg();
5197 Register CoordinateReg =
I.getOperand(4).getReg();
5198 ImageOperands ImOps;
5199 ImOps.Bias =
I.getOperand(5).getReg();
5200 if (
I.getNumOperands() > 6)
5201 ImOps.Offset =
I.getOperand(6).getReg();
5202 if (
I.getNumOperands() > 7)
5203 ImOps.MinLod =
I.getOperand(7).getReg();
5204 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5205 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5208bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
5209 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5210 Register ImageReg =
I.getOperand(2).getReg();
5211 Register SamplerReg =
I.getOperand(3).getReg();
5212 Register CoordinateReg =
I.getOperand(4).getReg();
5213 ImageOperands ImOps;
5214 ImOps.GradX =
I.getOperand(5).getReg();
5215 ImOps.GradY =
I.getOperand(6).getReg();
5216 if (
I.getNumOperands() > 7)
5217 ImOps.Offset =
I.getOperand(7).getReg();
5218 if (
I.getNumOperands() > 8)
5219 ImOps.MinLod =
I.getOperand(8).getReg();
5220 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5221 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5224bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
5225 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5226 Register ImageReg =
I.getOperand(2).getReg();
5227 Register SamplerReg =
I.getOperand(3).getReg();
5228 Register CoordinateReg =
I.getOperand(4).getReg();
5229 ImageOperands ImOps;
5230 ImOps.Lod =
I.getOperand(5).getReg();
5231 if (
I.getNumOperands() > 6)
5232 ImOps.Offset =
I.getOperand(6).getReg();
5233 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5234 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5237bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
5238 SPIRVTypeInst ResType,
5239 MachineInstr &
I)
const {
5240 Register ImageReg =
I.getOperand(2).getReg();
5241 Register SamplerReg =
I.getOperand(3).getReg();
5242 Register CoordinateReg =
I.getOperand(4).getReg();
5243 ImageOperands ImOps;
5244 ImOps.Compare =
I.getOperand(5).getReg();
5245 if (
I.getNumOperands() > 6)
5246 ImOps.Offset =
I.getOperand(6).getReg();
5247 if (
I.getNumOperands() > 7)
5248 ImOps.MinLod =
I.getOperand(7).getReg();
5249 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5250 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5253bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
5254 SPIRVTypeInst ResType,
5255 MachineInstr &
I)
const {
5256 Register ImageReg =
I.getOperand(2).getReg();
5257 Register CoordinateReg =
I.getOperand(3).getReg();
5258 Register LodReg =
I.getOperand(4).getReg();
5260 ImageOperands ImOps;
5262 if (
I.getNumOperands() > 5)
5263 ImOps.Offset =
I.getOperand(5).getReg();
5275 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
5276 I.getDebugLoc(),
I, &ImOps);
5279bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
5280 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5281 Register ImageReg =
I.getOperand(2).getReg();
5282 Register SamplerReg =
I.getOperand(3).getReg();
5283 Register CoordinateReg =
I.getOperand(4).getReg();
5284 ImageOperands ImOps;
5285 ImOps.Compare =
I.getOperand(5).getReg();
5286 if (
I.getNumOperands() > 6)
5287 ImOps.Offset =
I.getOperand(6).getReg();
5290 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5291 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5294bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
5295 SPIRVTypeInst ResType,
5296 MachineInstr &
I)
const {
5297 Register ImageReg =
I.getOperand(2).getReg();
5298 Register SamplerReg =
I.getOperand(3).getReg();
5299 Register CoordinateReg =
I.getOperand(4).getReg();
5302 "ImageReg is not an image type.");
5307 ComponentOrCompareReg =
I.getOperand(5).getReg();
5308 OffsetReg =
I.getOperand(6).getReg();
5311 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
5315 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5316 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
5317 Dim != SPIRV::Dim::DIM_Rect) {
5319 "Gather operations are only supported for 2D, Cube, and Rect images.");
5326 if (!loadHandleBeforePosition(
5331 MachineIRBuilder MIRBuilder(
I);
5332 SPIRVTypeInst SampledImageType =
5337 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5345 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
5347 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
5349 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5354 .
addUse(ComponentOrCompareReg);
5356 uint32_t ImageOperands = 0;
5357 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
5358 if (Dim == SPIRV::Dim::DIM_Cube) {
5360 "Gather operations with offset are not supported for Cube images.");
5364 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5366 ImageOperands |= SPIRV::ImageOperand::Offset;
5370 if (ImageOperands != 0) {
5371 MIB.
addImm(ImageOperands);
5373 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5381bool SPIRVInstructionSelector::generateImageReadOrFetch(
5384 const ImageOperands *ImOps)
const {
5387 "ImageReg is not an image type.");
5389 bool IsSignedInteger =
5394 bool IsFetch = (SampledOp.getImm() == 1);
5396 auto AddOperands = [&](MachineInstrBuilder &MIB) {
5397 uint32_t ImageOperandsMask = 0;
5398 if (IsSignedInteger)
5399 ImageOperandsMask |= 0x1000;
5401 if (IsFetch && ImOps) {
5403 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
5404 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
5406 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
5408 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5412 if (ImageOperandsMask != 0) {
5413 MIB.
addImm(ImageOperandsMask);
5414 if (IsFetch && ImOps) {
5417 if (ImOps->Offset &&
5418 (ImageOperandsMask &
5419 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5420 MIB.
addUse(*ImOps->Offset);
5426 if (ResultSize == 4) {
5429 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5436 BMI.constrainAllUses(
TII,
TRI, RBI);
5440 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5444 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5450 BMI.constrainAllUses(
TII,
TRI, RBI);
5452 if (ResultSize == 1) {
5461 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
5464bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
5465 SPIRVTypeInst ResType,
5466 MachineInstr &
I)
const {
5467 Register ResourcePtr =
I.getOperand(2).getReg();
5469 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
5478 MachineIRBuilder MIRBuilder(
I);
5480 Register IndexReg =
I.getOperand(3).getReg();
5483 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5493bool SPIRVInstructionSelector::selectPushConstantGetPointer(
5494 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5499bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
5500 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5501 Register ObjReg =
I.getOperand(2).getReg();
5502 if (!BuildCOPY(ResVReg, ObjReg,
I))
5512 decorateUsesAsNonUniform(ResVReg);
5516void SPIRVInstructionSelector::decorateUsesAsNonUniform(
5519 while (WorkList.
size() > 0) {
5523 bool IsDecorated =
false;
5525 if (
Use.getOpcode() == SPIRV::OpDecorate &&
5526 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
5532 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
5534 if (ResultReg == CurrentReg)
5542 SPIRV::Decoration::NonUniformEXT, {});
5547bool SPIRVInstructionSelector::extractSubvector(
5549 MachineInstr &InsertionPoint)
const {
5551 [[maybe_unused]] uint64_t InputSize =
5554 assert(InputSize > 1 &&
"The input must be a vector.");
5555 assert(ResultSize > 1 &&
"The result must be a vector.");
5556 assert(ResultSize < InputSize &&
5557 "Cannot extract more element than there are in the input.");
5560 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
5561 for (uint64_t
I = 0;
I < ResultSize;
I++) {
5564 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5573 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
5575 TII.get(SPIRV::OpCompositeConstruct))
5579 for (
Register ComponentReg : ComponentRegisters)
5580 MIB.
addUse(ComponentReg);
5585bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
5586 MachineInstr &
I)
const {
5593 Register ImageReg =
I.getOperand(1).getReg();
5601 Register CoordinateReg =
I.getOperand(2).getReg();
5602 Register DataReg =
I.getOperand(3).getReg();
5605 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
5613Register SPIRVInstructionSelector::buildPointerToResource(
5614 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
5615 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
5616 StringRef Name, MachineIRBuilder MIRBuilder)
const {
5618 if (ArraySize == 1) {
5619 SPIRVTypeInst PtrType =
5622 "SpirvResType did not have an explicit layout.");
5627 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
5628 SPIRVTypeInst VarPointerType =
5631 VarPointerType, Set,
Binding, Name, MIRBuilder);
5633 SPIRVTypeInst ResPointerType =
5646bool SPIRVInstructionSelector::selectFirstBitSet16(
5647 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5648 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
5650 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
5654 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
5657bool SPIRVInstructionSelector::selectFirstBitSet32(
5659 unsigned BitSetOpcode)
const {
5660 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5663 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5670bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
5672 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5679 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
5681 MachineIRBuilder MIRBuilder(
I);
5684 SPIRVTypeInst I64x2Type =
5686 SPIRVTypeInst Vec2ResType =
5689 std::vector<Register> PartialRegs;
5692 unsigned CurrentComponent = 0;
5693 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
5699 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5700 TII.get(SPIRV::OpVectorShuffle))
5705 .
addImm(CurrentComponent)
5706 .
addImm(CurrentComponent + 1);
5713 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType,
I, BitSetResult,
5714 BitSetOpcode, SwapPrimarySide))
5717 PartialRegs.push_back(SubVecBitSetReg);
5721 if (CurrentComponent != ComponentCount) {
5727 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
5728 SPIRV::OpVectorExtractDynamic))
5734 if (!selectFirstBitSet64(FinalElemBitSetReg,
BaseType,
I, FinalElemReg,
5735 BitSetOpcode, SwapPrimarySide))
5738 PartialRegs.push_back(FinalElemBitSetReg);
5743 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
5744 SPIRV::OpCompositeConstruct);
5747bool SPIRVInstructionSelector::selectFirstBitSet64(
5749 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5762 if (ComponentCount > 2) {
5763 return selectFirstBitSet64Overflow(ResVReg, ResType,
I, SrcReg,
5764 BitSetOpcode, SwapPrimarySide);
5768 MachineIRBuilder MIRBuilder(
I);
5770 BaseType, 2 * ComponentCount, MIRBuilder,
false);
5774 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
5780 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
5787 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
5790 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
5791 SPIRV::OpVectorExtractDynamic))
5793 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
5794 SPIRV::OpVectorExtractDynamic))
5798 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5799 TII.get(SPIRV::OpVectorShuffle))
5807 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
5813 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5814 TII.get(SPIRV::OpVectorShuffle))
5822 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
5842 SelectOp = SPIRV::OpSelectSISCond;
5843 AddOp = SPIRV::OpIAddS;
5851 SelectOp = SPIRV::OpSelectVIVCond;
5852 AddOp = SPIRV::OpIAddV;
5858 Register RegSecondaryOffset = Reg0;
5862 if (SwapPrimarySide) {
5863 PrimaryReg = LowReg;
5864 SecondaryReg = HighReg;
5865 RegPrimaryOffset = Reg0;
5866 RegSecondaryOffset = Reg32;
5871 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
5872 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
5877 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
5878 SPIRV::OpINotEqual))
5885 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
5886 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
5891 if (SwapPrimarySide) {
5893 if (!selectOpWithSrcs(RegAdd, ResType,
I,
5894 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
5905 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
5906 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
5911 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
5912 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
5915 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
5919bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
5920 SPIRVTypeInst ResType,
5922 bool IsSigned)
const {
5924 Register OpReg =
I.getOperand(2).getReg();
5927 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5928 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
5932 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5934 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5936 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5940 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
5944bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
5945 SPIRVTypeInst ResType,
5946 MachineInstr &
I)
const {
5948 Register OpReg =
I.getOperand(2).getReg();
5953 unsigned ExtendOpcode = SPIRV::OpUConvert;
5954 unsigned BitSetOpcode = GL::FindILsb;
5958 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5960 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5962 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5969bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
5970 SPIRVTypeInst ResType,
5971 MachineInstr &
I)
const {
5975 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
5978 .
addUse(
I.getOperand(2).getReg())
5981 unsigned Alignment =
I.getOperand(3).getImm();
5987bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
5988 SPIRVTypeInst ResType,
5989 MachineInstr &
I)
const {
5993 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
5996 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
5999 unsigned Alignment =
I.getOperand(2).getImm();
6006bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6011 const MachineInstr *PrevI =
I.getPrevNode();
6013 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6017 .
addMBB(
I.getOperand(0).getMBB())
6022 .
addMBB(
I.getOperand(0).getMBB())
6027bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6038 const MachineInstr *NextI =
I.getNextNode();
6040 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6046 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6048 .
addUse(
I.getOperand(0).getReg())
6049 .
addMBB(
I.getOperand(1).getMBB())
6055bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6056 MachineInstr &
I)
const {
6058 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6060 const unsigned NumOps =
I.getNumOperands();
6061 for (
unsigned i = 1; i <
NumOps; i += 2) {
6062 MIB.
addUse(
I.getOperand(i + 0).getReg());
6063 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6069bool SPIRVInstructionSelector::selectGlobalValue(
6070 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6072 MachineIRBuilder MIRBuilder(
I);
6073 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6076 std::string GlobalIdent;
6078 unsigned &
ID = UnnamedGlobalIDs[GV];
6080 ID = UnnamedGlobalIDs.
size();
6081 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6107 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6114 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6119 MachineInstrBuilder MIB1 =
6120 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6123 MachineInstrBuilder MIB2 =
6125 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
6129 GR.
add(ConstVal, MIB2);
6137 MachineInstrBuilder MIB3 =
6138 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
6141 GR.
add(ConstVal, MIB3);
6145 assert(NewReg != ResVReg);
6146 return BuildCOPY(ResVReg, NewReg,
I);
6156 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
6159 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
6165 SPIRVTypeInst ResType =
6169 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6174 if (
GlobalVar->isExternallyInitialized() &&
6175 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6176 constexpr unsigned ReadWriteINTEL = 3u;
6179 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6185bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6186 SPIRVTypeInst ResType,
6187 MachineInstr &
I)
const {
6189 return selectExtInst(ResVReg, ResType,
I, CL::log10);
6197 MachineIRBuilder MIRBuilder(
I);
6202 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6205 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6207 .
add(
I.getOperand(1))
6212 ResType->
getOpcode() == SPIRV::OpTypeFloat);
6214 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
6222 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
6223 ? SPIRV::OpVectorTimesScalar
6234bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
6235 SPIRVTypeInst ResType,
6236 MachineInstr &
I)
const {
6239 return selectExtInst(ResVReg, ResType,
I, CL::pown);
6245 Register ExpReg =
I.getOperand(2).getReg();
6247 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
6248 SPIRV::OpConvertSToF))
6250 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
6257bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
6258 SPIRVTypeInst ResType,
6259 MachineInstr &
I)
const {
6275 MachineIRBuilder MIRBuilder(
I);
6278 ResType, MIRBuilder, SPIRV::StorageClass::Function);
6290 MachineBasicBlock &EntryBB =
I.getMF()->front();
6294 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
6297 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
6303 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6306 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
6309 .
add(
I.getOperand(
I.getNumExplicitDefs()))
6313 Register IntegralPartReg =
I.getOperand(1).getReg();
6314 if (IntegralPartReg.
isValid()) {
6316 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6327 assert(
false &&
"GLSL::Modf is deprecated.");
6338bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
6339 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6340 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6341 MachineIRBuilder MIRBuilder(
I);
6342 const SPIRVTypeInst Vec3Ty =
6345 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
6357 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6361 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6367 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6374 assert(
I.getOperand(2).isReg());
6375 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
6379 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6390bool SPIRVInstructionSelector::loadBuiltinInputID(
6391 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6392 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6393 MachineIRBuilder MIRBuilder(
I);
6395 ResType, MIRBuilder, SPIRV::StorageClass::Input);
6410 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6414 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6423SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
6424 MachineInstr &
I)
const {
6425 MachineIRBuilder MIRBuilder(
I);
6426 if (
Type->getOpcode() != SPIRV::OpTypeVector)
6430 if (VectorSize == 4)
6438bool SPIRVInstructionSelector::loadHandleBeforePosition(
6439 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
6440 MachineInstr &Pos)
const {
6443 Intrinsic::spv_resource_handlefrombinding);
6451 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
6452 MachineIRBuilder MIRBuilder(HandleDef);
6453 SPIRVTypeInst VarType = ResType;
6454 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
6456 if (IsStructuredBuffer) {
6461 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
6463 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
6466 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
6467 ArraySize, IndexReg, Name, MIRBuilder);
6471 uint32_t LoadOpcode =
6472 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
6482void SPIRVInstructionSelector::errorIfInstrOutsideShader(
6483 MachineInstr &
I)
const {
6485 std::string DiagMsg;
6486 raw_string_ostream OS(DiagMsg);
6487 I.print(OS,
true,
false,
false,
false);
6488 DiagMsg +=
" is only supported in shaders.\n";
6494InstructionSelector *
6498 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
MachineInstrBuilder & UseMI
#define GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
This file declares a class to represent arbitrary precision floating point values and provide a varie...
static bool selectUnmergeValues(MachineInstrBuilder &MIB, const ARMBaseInstrInfo &TII, MachineRegisterInfo &MRI, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
static uint8_t SwapBits(uint8_t Val)
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
DXIL Resource Implicit Binding
Declares convenience wrapper classes for interpreting MachineInstr instances as specific generic oper...
const HexagonInstrInfo * TII
LLVMTypeRef LLVMIntType(unsigned NumBits)
const size_t AbstractManglingParser< Derived, Alloc >::NumOps
Loop::LoopBounds::Direction Direction
Register const TargetRegisterInfo * TRI
Promote Memory to Register
MachineInstr unsigned OpIdx
uint64_t IntrinsicInst * II
static StringRef getName(Value *V)
static unsigned getFCmpOpcode(CmpInst::Predicate Pred, unsigned Size)
static APFloat getOneFP(const Type *LLVMFloatTy)
static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC)
static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg)
static bool mayApplyGenericSelection(unsigned Opcode)
static APFloat getZeroFP(const Type *LLVMFloatTy)
std::vector< std::pair< SPIRV::InstructionSet::InstructionSet, uint32_t > > ExtInstList
static bool intrinsicHasSideEffects(Intrinsic::ID ID)
static unsigned getBoolCmpOpcode(unsigned PredNum)
static unsigned getICmpOpcode(unsigned PredNum)
static bool isOpcodeWithNoSideEffects(unsigned Opcode)
static void addMemoryOperands(MachineMemOperand *MemOp, MachineInstrBuilder &MIB, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry &GR)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef)
static unsigned getPtrCmpOpcode(unsigned Pred)
unsigned getVectorSizeOrOne(SPIRVTypeInst Type)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")
static const fltSemantics & IEEEsingle()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
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.
ArrayRef - Represent a constant reference to an array (0 or more elements consecutively in memory),...
BlockFrequencyInfo pass uses BlockFrequencyInfoImpl implementation to estimate IR basic block frequen...
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
Represents a call to an intrinsic.
Intrinsic::ID getIntrinsicID() const
unsigned getAddressSpace() const
Module * getParent()
Get the module that this global value is contained inside of...
@ InternalLinkage
Rename collisions when linking (static functions).
static LLVM_ABI IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
constexpr bool isValid() const
constexpr uint16_t getNumElements() const
Returns the number of elements in a vector LLT.
constexpr bool isVector() const
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
int getNumber() const
MachineBasicBlocks are uniquely numbered at the function level, unless they're not in a MachineFuncti...
LLVM_ABI iterator getFirstNonPHI()
Returns a pointer to the first instruction in this block that is not a PHINode instruction.
const MachineFunction * getParent() const
Return the MachineFunction containing this basic block.
MachineInstrBundleIterator< MachineInstr > iterator
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
Helper class to build MachineInstr.
MachineBasicBlock::iterator getInsertPt()
Current insertion point for new instructions.
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
MachineFunction & getMF()
Getter for the function we currently build.
MachineRegisterInfo * getMRI()
Getter for MRI.
void constrainAllUses(const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI) const
const MachineInstrBuilder & addUse(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register use operand.
const MachineInstrBuilder & addReg(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a new virtual register operand.
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
const MachineInstrBuilder & add(const MachineOperand &MO) const
const MachineInstrBuilder & addMBB(MachineBasicBlock *MBB, unsigned TargetFlags=0) const
const MachineInstrBuilder & addDef(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register definition operand.
const MachineInstrBuilder & setMIFlags(unsigned Flags) const
MachineInstr * getInstr() const
If conversion operators fail, use this method to get the MachineInstr explicitly.
Representation of each machine instruction.
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
const MachineBasicBlock * getParent() const
unsigned getNumOperands() const
Retuns the total number of operands.
LLVM_ABI unsigned getNumExplicitOperands() const
Returns the number of non-implicit operands.
LLVM_ABI unsigned getNumExplicitDefs() const
Returns the number of non-implicit definitions.
LLVM_ABI void emitGenericError(const Twine &ErrMsg) const
LLVM_ABI const MachineFunction * getMF() const
Return the function that contains the basic block that this instruction belongs to.
const DebugLoc & getDebugLoc() const
Returns the debug location id of this MachineInstr.
const MachineOperand & getOperand(unsigned i) const
A description of a memory reference used in the backend.
@ MOVolatile
The memory access is volatile.
@ MONonTemporal
The memory access is non-temporal.
bool isReg() const
isReg - Tests if this is a MO_Register operand.
MachineBasicBlock * getMBB() const
Register getReg() const
getReg - Returns the register number.
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
defusechain_instr_iterator< true, false, false, true > use_instr_iterator
use_instr_iterator/use_instr_begin/use_instr_end - Walk all uses of the specified register,...
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
LLVM_ABI MachineInstr * getVRegDef(Register Reg) const
getVRegDef - Return the machine instr that defines the specified virtual register or null if none is ...
use_instr_iterator use_instr_begin(Register RegNo) const
static def_instr_iterator def_instr_end()
defusechain_instr_iterator< false, true, false, true > def_instr_iterator
def_instr_iterator/def_instr_begin/def_instr_end - Walk all defs of the specified register,...
LLVM_ABI Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
def_instr_iterator def_instr_begin(Register RegNo) const
LLT getType(Register Reg) const
Get the low-level type of Reg or LLT{} if Reg is not a generic (target independent) virtual register.
static use_instr_iterator use_instr_end()
iterator_range< use_instr_nodbg_iterator > use_nodbg_instructions(Register Reg) const
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
const MachineFunction & getMF() const
LLVM_ABI void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
LLVM_ABI Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
const TargetRegisterClass * getRegClassOrNull(Register Reg) const
Return the register class of Reg, or null if Reg has not been assigned a register class yet.
iterator_range< use_instr_iterator > use_instructions(Register Reg) const
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
LLVM_ABI void replaceRegWith(Register FromReg, Register ToReg)
replaceRegWith - Replace all instances of FromReg with ToReg in the machine function.
Analysis providing profile information.
Holds all the information related to register banks.
Wrapper class representing virtual and physical registers.
constexpr bool isValid() const
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
bool isScalarOrVectorSigned(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateOpTypeSampledImage(SPIRVTypeInst ImageType, MachineIRBuilder &MIRBuilder)
void assignSPIRVTypeToVReg(SPIRVTypeInst Type, Register VReg, const MachineFunction &MF)
const TargetRegisterClass * getRegClass(SPIRVTypeInst SpvType) const
MachineInstr * getOrAddMemAliasingINTELInst(MachineIRBuilder &MIRBuilder, const MDNode *AliasingListMD)
bool isAggregateType(SPIRVTypeInst Type) const
unsigned getScalarOrVectorBitWidth(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst getOrCreateSPIRVVectorType(SPIRVTypeInst BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
Register buildGlobalVariable(Register Reg, SPIRVTypeInst BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
SPIRVTypeInst getResultType(Register VReg, MachineFunction *MF=nullptr)
unsigned getScalarOrVectorComponentCount(Register VReg) const
const Type * getTypeForSPIRVType(SPIRVTypeInst Ty) const
bool isBitcastCompatible(SPIRVTypeInst Type1, SPIRVTypeInst Type2) const
unsigned getPointerSize() const
Register getOrCreateConstFP(APFloat Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
LLT getRegType(SPIRVTypeInst SpvType) const
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC)
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
Register getSPIRVTypeID(SPIRVTypeInst SpirvType) const
Register getOrCreateConstInt(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
bool findValueAttrs(const MachineInstr *Key, Type *&Ty, StringRef &Name)
SPIRVTypeInst retrieveScalarOrVectorIntType(SPIRVTypeInst Type) const
Register getOrCreateGlobalVariableWithBinding(SPIRVTypeInst VarType, uint32_t Set, uint32_t Binding, StringRef Name, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst changePointerStorageClass(SPIRVTypeInst PtrType, SPIRV::StorageClass::StorageClass SC, MachineInstr &I)
Register getOrCreateConstVector(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType=nullptr)
void addGlobalObject(const Value *V, const MachineFunction *MF, Register R)
SPIRVTypeInst getScalarOrVectorComponentType(SPIRVTypeInst Type) const
void recordFunctionPointer(const MachineOperand *MO, const Function *F)
SPIRVTypeInst getOrCreateSPIRVFloatType(unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII)
SPIRVTypeInst getPointeeType(SPIRVTypeInst PtrType)
SPIRVTypeInst getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
MachineFunction * setCurrentFunc(MachineFunction &MF)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType)
SPIRVTypeInst getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
Type * getDeducedGlobalValueType(const GlobalValue *Global)
Register getOrCreateUndef(MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
bool erase(const MachineInstr *MI)
bool add(SPIRV::IRHandle Handle, const MachineInstr *MI)
Register find(SPIRV::IRHandle Handle, const MachineFunction *MF)
bool isPhysicalSPIRV() const
bool isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const
bool canUseExtInstSet(SPIRV::InstructionSet::InstructionSet E) const
bool isLogicalSPIRV() const
bool canUseExtension(SPIRV::Extension::Extension E) const
bool erase(PtrType Ptr)
Remove pointer from the set.
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
bool contains(ConstPtrType Ptr) const
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
reference emplace_back(ArgTypes &&... Args)
void reserve(size_type N)
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
constexpr size_t size() const
size - 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
@ DoubleTyID
64-bit floating point type
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
bool isStructTy() const
True if this is an instance of StructType.
bool isAggregateType() const
Return true if the type is an aggregate type.
TypeID getTypeID() const
Return the type id for the type.
Value * getOperand(unsigned i) const
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
self_iterator getIterator()
NodeTy * getNextNode()
Get the next node, or nullptr for the list tail.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
constexpr char IsConst[]
Key for Kernel::Arg::Metadata::mIsConst.
constexpr std::underlying_type_t< E > Mask()
Get a bitmask with 1s in all places up to the high-order bit of E's largest value.
unsigned ID
LLVM IR allows to use arbitrary numbers as calling convention identifiers.
Scope
Defines the scope in which this symbol should be visible: Default – Visible in the public interface o...
NodeAddr< DefNode * > Def
NodeAddr< InstrNode * > Instr
NodeAddr< UseNode * > Use
BaseReg
Stack frame base register. Bit 0 of FREInfo.Info.
This is an optimization pass for GlobalISel generic memory operations.
void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)
FunctionAddr VTableAddr Value
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
int64_t getIConstValSext(Register ConstReg, const MachineRegisterInfo *MRI)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
FunctionAddr VTableAddr uintptr_t uintptr_t Int32Ty
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
LLVM_ABI void salvageDebugInfo(const MachineRegisterInfo &MRI, MachineInstr &MI)
Assuming the instruction MI is going to be deleted, attempt to salvage debug users of MI by writing t...
LLVM_ABI void constrainSelectedInstRegOperands(MachineInstr &I, const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
Mutate the newly-selected instruction I to constrain its (possibly generic) virtual register operands...
bool isPreISelGenericOpcode(unsigned Opcode)
Check whether the given Opcode is a generic opcode that is not supposed to appear after ISel.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
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)
MachineBasicBlock::iterator getFirstValidInstructionInsertPoint(MachineBasicBlock &BB)
bool any_of(R &&range, UnaryPredicate P)
Provide wrappers to std::any_of which take ranges instead of having to pass begin/end explicitly.
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineInstr &I)
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
Type * toTypedPointer(Type *Ty)
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
constexpr bool isGenericCastablePtr(SPIRV::StorageClass::StorageClass SC)
class LLVM_GSL_OWNER SmallVector
Forward declaration of SmallVector so that calculateSmallVectorDefaultInlinedElements can reference s...
MachineInstr * passCopy(MachineInstr *Def, const MachineRegisterInfo *MRI)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
std::optional< SPIRV::LinkageType::LinkageType > getSpirvLinkageTypeFor(const SPIRVSubtarget &ST, const GlobalValue &GV)
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
SPIRV::StorageClass::StorageClass addressSpaceToStorageClass(unsigned AddrSpace, const SPIRVSubtarget &STI)
AtomicOrdering
Atomic ordering for LLVM's memory model.
SPIRV::Scope::Scope getMemScope(LLVMContext &Ctx, SyncScope::ID Id)
InstructionSelector * createSPIRVInstructionSelector(const SPIRVTargetMachine &TM, const SPIRVSubtarget &Subtarget, const RegisterBankInfo &RBI)
std::string getStringValueFromReg(Register Reg, MachineRegisterInfo &MRI)
int64_t foldImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
DWARFExpression::Operation Op
ArrayRef(const T &OneElt) -> ArrayRef< T >
MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)
constexpr unsigned BitWidth
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
bool hasInitializer(const GlobalVariable *GV)
bool isSpvIntrinsic(const MachineInstr &MI, Intrinsic::ID IntrinsicID)
void addStringImm(const StringRef &Str, MCInst &Inst)
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...