Lines Matching +full:dim +full:- +full:mode
1 //===- SPIRVBuiltins.cpp - SPIR-V Built-in Functions ------------*- C++ -*-===//
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7 //===----------------------------------------------------------------------===//
12 //===----------------------------------------------------------------------===//
24 #define DEBUG_TYPE "spirv-builtins"
167 //===----------------------------------------------------------------------===//
170 //===----------------------------------------------------------------------===//
186 // Account for possible "__spirv_ocl_" prefix in SPIR-V friendly LLVM IR in lookupBuiltinNameHelper()
191 // brackets. If so, the builtin is an instantiated template - needs to have in lookupBuiltinNameHelper()
230 BuiltinArgs.split(BuiltinArgumentTypes, ',', -1, false); in lookupBuiltin()
266 // Floating-point: in lookupBuiltin()
276 // If argument-type name prefix was added, look up the builtin again. in lookupBuiltin()
299 // Floating-point: in lookupBuiltin()
307 // If argument-type name suffix was added, look up the builtin again. in lookupBuiltin()
325 MachineInstr *MI = MRI->getUniqueVRegDef(ParamReg); in getBlockStructInstr()
326 assert(MI->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST && in getBlockStructInstr()
327 MI->getOperand(1).isReg()); in getBlockStructInstr()
328 Register BitcastReg = MI->getOperand(1).getReg(); in getBlockStructInstr()
329 MachineInstr *BitcastMI = MRI->getUniqueVRegDef(BitcastReg); in getBlockStructInstr()
331 BitcastMI->getOperand(2).isReg()); in getBlockStructInstr()
332 Register ValueReg = BitcastMI->getOperand(2).getReg(); in getBlockStructInstr()
333 MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg); in getBlockStructInstr()
341 MachineInstr *DefMI = MRI->getUniqueVRegDef(Reg); in getConstFromIntrinsic()
343 DefMI->getOperand(2).isReg()); in getConstFromIntrinsic()
344 MachineInstr *DefMI2 = MRI->getUniqueVRegDef(DefMI->getOperand(2).getReg()); in getConstFromIntrinsic()
345 assert(DefMI2->getOpcode() == TargetOpcode::G_CONSTANT && in getConstFromIntrinsic()
346 DefMI2->getOperand(1).isCImm()); in getConstFromIntrinsic()
347 return DefMI2->getOperand(1).getCImm()->getValue().getZExtValue(); in getConstFromIntrinsic()
353 MachineInstr *NextMI = MI->getNextNode(); in getMachineInstrType()
357 if ((NextMI = NextMI->getNextNode()) == nullptr) in getMachineInstrType()
359 Register ValueReg = MI->getOperand(0).getReg(); in getMachineInstrType()
362 NextMI->getOperand(1).getReg() != ValueReg) in getMachineInstrType()
364 Type *Ty = getMDOperandAsType(NextMI->getOperand(2).getMetadata(), 0); in getMachineInstrType()
377 if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) in getBlockStructType()
378 return MI->getOperand(1).getGlobal()->getType(); in getBlockStructType()
384 //===----------------------------------------------------------------------===//
386 //===----------------------------------------------------------------------===//
396 SPIRVType *BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder); in buildBoolRegister()
398 if (ResultType->getOpcode() == SPIRV::OpTypeVector) { in buildBoolRegister()
399 unsigned VectorElements = ResultType->getOperand(2).getImm(); in buildBoolRegister()
401 GR->getOrCreateSPIRVVectorType(BoolType, VectorElements, MIRBuilder); in buildBoolRegister()
403 cast<FixedVectorType>(GR->getTypeForSPIRVType(BoolType)); in buildBoolRegister()
404 Type = LLT::vector(LLVMVectorType->getElementCount(), 1); in buildBoolRegister()
410 MIRBuilder.getMRI()->createGenericVirtualRegister(Type); in buildBoolRegister()
411 MIRBuilder.getMRI()->setRegClass(ResultRegister, &SPIRV::IDRegClass); in buildBoolRegister()
412 GR->assignSPIRVTypeToVReg(BoolType, ResultRegister, MIRBuilder.getMF()); in buildBoolRegister()
424 if (ReturnType->getOpcode() == SPIRV::OpTypeVector) { in buildSelectInst()
425 unsigned Bits = GR->getScalarOrVectorBitWidth(ReturnType); in buildSelectInst()
427 TrueConst = GR->getOrCreateConsIntVector(AllOnes, MIRBuilder, ReturnType); in buildSelectInst()
428 FalseConst = GR->getOrCreateConsIntVector(0, MIRBuilder, ReturnType); in buildSelectInst()
430 TrueConst = GR->buildConstantInt(1, MIRBuilder, ReturnType); in buildSelectInst()
431 FalseConst = GR->buildConstantInt(0, MIRBuilder, ReturnType); in buildSelectInst()
445 DestinationReg = MRI->createVirtualRegister(&SPIRV::IDRegClass); in buildLoadInst()
446 MRI->setType(DestinationReg, LLT::scalar(32)); in buildLoadInst()
447 GR->assignSPIRVTypeToVReg(BaseType, DestinationReg, MIRBuilder.getMF()); in buildLoadInst()
463 MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass); in buildBuiltinVariableLoad()
464 MIRBuilder.getMRI()->setType(NewRegister, in buildBuiltinVariableLoad()
465 LLT::pointer(0, GR->getPointerSize())); in buildBuiltinVariableLoad()
466 SPIRVType *PtrType = GR->getOrCreateSPIRVPointerType( in buildBuiltinVariableLoad()
468 GR->assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF()); in buildBuiltinVariableLoad()
471 Register Variable = GR->buildGlobalVariable( in buildBuiltinVariableLoad()
480 MIRBuilder.getMRI()->setType(LoadedRegister, LLType); in buildBuiltinVariableLoad()
532 SPIRVType *IntType = GR->getOrCreateSPIRVIntegerType(BitWidth, MIRBuilder); in buildConstantIntReg()
533 return GR->buildConstantInt(Val, MIRBuilder, IntType); in buildConstantIntReg()
547 MRI->setRegClass(CLScopeRegister, &SPIRV::IDRegClass); in buildScopeReg()
564 getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister)); in buildMemSemanticsReg()
567 MRI->setRegClass(SemanticsRegister, &SPIRV::IDRegClass); in buildMemSemanticsReg()
581 MIB.addDef(Call->ReturnRegister).addUse(TypeReg);
582 unsigned Sz = Call->Arguments.size() - ImmArgs.size();
584 Register ArgReg = Call->Arguments[i];
585 if (!MRI->getRegClassOrNull(ArgReg))
586 MRI->setRegClass(ArgReg, &SPIRV::IDRegClass);
597 if (Call->isSpirvOp()) in buildAtomicInitInst()
600 assert(Call->Arguments.size() == 2 && in buildAtomicInitInst()
602 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass); in buildAtomicInitInst()
603 MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass); in buildAtomicInitInst()
605 .addUse(Call->Arguments[0]) in buildAtomicInitInst()
606 .addUse(Call->Arguments[1]); in buildAtomicInitInst()
614 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType); in buildAtomicLoadInst()
615 if (Call->isSpirvOp()) in buildAtomicLoadInst()
618 Register PtrRegister = Call->Arguments[0]; in buildAtomicLoadInst()
619 MIRBuilder.getMRI()->setRegClass(PtrRegister, &SPIRV::IDRegClass); in buildAtomicLoadInst()
624 if (Call->Arguments.size() > 1) { in buildAtomicLoadInst()
625 ScopeRegister = Call->Arguments[1]; in buildAtomicLoadInst()
626 MIRBuilder.getMRI()->setRegClass(ScopeRegister, &SPIRV::IDRegClass); in buildAtomicLoadInst()
631 if (Call->Arguments.size() > 2) { in buildAtomicLoadInst()
633 MemSemanticsReg = Call->Arguments[2]; in buildAtomicLoadInst()
634 MIRBuilder.getMRI()->setRegClass(MemSemanticsReg, &SPIRV::IDRegClass); in buildAtomicLoadInst()
638 getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister)); in buildAtomicLoadInst()
643 .addDef(Call->ReturnRegister) in buildAtomicLoadInst()
655 if (Call->isSpirvOp()) in buildAtomicStoreInst()
660 Register PtrRegister = Call->Arguments[0]; in buildAtomicStoreInst()
661 MIRBuilder.getMRI()->setRegClass(PtrRegister, &SPIRV::IDRegClass); in buildAtomicStoreInst()
664 getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister)); in buildAtomicStoreInst()
666 MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass); in buildAtomicStoreInst()
671 .addUse(Call->Arguments[1]); in buildAtomicStoreInst()
675 /// Helper function for building an atomic compare-exchange instruction.
679 if (Call->isSpirvOp()) in buildAtomicCompareExchangeInst()
681 GR->getSPIRVTypeID(Call->ReturnType)); in buildAtomicCompareExchangeInst()
683 bool IsCmpxchg = Call->Builtin->Name.contains("cmpxchg"); in buildAtomicCompareExchangeInst()
686 Register ObjectPtr = Call->Arguments[0]; // Pointer (volatile A *object.) in buildAtomicCompareExchangeInst()
687 Register ExpectedArg = Call->Arguments[1]; // Comparator (C* expected). in buildAtomicCompareExchangeInst()
688 Register Desired = Call->Arguments[2]; // Value (C Desired). in buildAtomicCompareExchangeInst()
689 MRI->setRegClass(ObjectPtr, &SPIRV::IDRegClass); in buildAtomicCompareExchangeInst()
690 MRI->setRegClass(ExpectedArg, &SPIRV::IDRegClass); in buildAtomicCompareExchangeInst()
691 MRI->setRegClass(Desired, &SPIRV::IDRegClass); in buildAtomicCompareExchangeInst()
692 SPIRVType *SpvDesiredTy = GR->getSPIRVTypeForVReg(Desired); in buildAtomicCompareExchangeInst()
693 LLT DesiredLLT = MRI->getType(Desired); in buildAtomicCompareExchangeInst()
695 assert(GR->getSPIRVTypeForVReg(ObjectPtr)->getOpcode() == in buildAtomicCompareExchangeInst()
697 unsigned ExpectedType = GR->getSPIRVTypeForVReg(ExpectedArg)->getOpcode(); in buildAtomicCompareExchangeInst()
701 assert(GR->isScalarOfType(Desired, SPIRV::OpTypeInt)); in buildAtomicCompareExchangeInst()
703 SPIRVType *SpvObjectPtrTy = GR->getSPIRVTypeForVReg(ObjectPtr); in buildAtomicCompareExchangeInst()
704 assert(SpvObjectPtrTy->getOperand(2).isReg() && "SPIRV type is expected"); in buildAtomicCompareExchangeInst()
706 SpvObjectPtrTy->getOperand(1).getImm()); in buildAtomicCompareExchangeInst()
719 if (Call->Arguments.size() >= 4) { in buildAtomicCompareExchangeInst()
720 assert(Call->Arguments.size() >= 5 && in buildAtomicCompareExchangeInst()
723 static_cast<std::memory_order>(getIConstVal(Call->Arguments[3], MRI)); in buildAtomicCompareExchangeInst()
725 static_cast<std::memory_order>(getIConstVal(Call->Arguments[4], MRI)); in buildAtomicCompareExchangeInst()
729 MemSemEqualReg = Call->Arguments[3]; in buildAtomicCompareExchangeInst()
731 MemSemUnequalReg = Call->Arguments[4]; in buildAtomicCompareExchangeInst()
732 MRI->setRegClass(Call->Arguments[3], &SPIRV::IDRegClass); in buildAtomicCompareExchangeInst()
733 MRI->setRegClass(Call->Arguments[4], &SPIRV::IDRegClass); in buildAtomicCompareExchangeInst()
742 if (Call->Arguments.size() >= 6) { in buildAtomicCompareExchangeInst()
743 assert(Call->Arguments.size() == 6 && in buildAtomicCompareExchangeInst()
746 getIConstVal(Call->Arguments[5], MRI)); in buildAtomicCompareExchangeInst()
749 ScopeReg = Call->Arguments[5]; in buildAtomicCompareExchangeInst()
750 MRI->setRegClass(Call->Arguments[5], &SPIRV::IDRegClass); in buildAtomicCompareExchangeInst()
759 MRI->setType(Expected, DesiredLLT); in buildAtomicCompareExchangeInst()
760 Register Tmp = !IsCmpxchg ? MRI->createGenericVirtualRegister(DesiredLLT) in buildAtomicCompareExchangeInst()
761 : Call->ReturnRegister; in buildAtomicCompareExchangeInst()
762 if (!MRI->getRegClassOrNull(Tmp)) in buildAtomicCompareExchangeInst()
763 MRI->setRegClass(Tmp, &SPIRV::IDRegClass); in buildAtomicCompareExchangeInst()
764 GR->assignSPIRVTypeToVReg(SpvDesiredTy, Tmp, MIRBuilder.getMF()); in buildAtomicCompareExchangeInst()
766 SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder); in buildAtomicCompareExchangeInst()
769 .addUse(GR->getSPIRVTypeID(IntTy)) in buildAtomicCompareExchangeInst()
778 MIRBuilder.buildICmp(CmpInst::ICMP_EQ, Call->ReturnRegister, Tmp, Expected); in buildAtomicCompareExchangeInst()
787 if (Call->isSpirvOp()) in buildAtomicRMWInst()
789 GR->getSPIRVTypeID(Call->ReturnType)); in buildAtomicRMWInst()
793 Call->Arguments.size() >= 4 ? Call->Arguments[3] : Register(); in buildAtomicRMWInst()
795 assert(Call->Arguments.size() <= 4 && in buildAtomicRMWInst()
800 Register PtrRegister = Call->Arguments[0]; in buildAtomicRMWInst()
802 MRI->setRegClass(PtrRegister, &SPIRV::IDRegClass); in buildAtomicRMWInst()
804 Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register(); in buildAtomicRMWInst()
807 MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass); in buildAtomicRMWInst()
808 Register ValueReg = Call->Arguments[1]; in buildAtomicRMWInst()
809 Register ValueTypeReg = GR->getSPIRVTypeID(Call->ReturnType); in buildAtomicRMWInst()
811 if (Call->ReturnType->getOpcode() == SPIRV::OpTypeFloat) { in buildAtomicRMWInst()
819 MRI->createGenericVirtualRegister(MRI->getType(ValueReg)); in buildAtomicRMWInst()
820 MRI->setRegClass(NegValueReg, &SPIRV::IDRegClass); in buildAtomicRMWInst()
821 GR->assignSPIRVTypeToVReg(Call->ReturnType, NegValueReg, in buildAtomicRMWInst()
826 insertAssignInstr(NegValueReg, nullptr, Call->ReturnType, GR, MIRBuilder, in buildAtomicRMWInst()
832 .addDef(Call->ReturnRegister) in buildAtomicRMWInst()
841 /// Helper function for building an atomic floating-type instruction.
846 assert(Call->Arguments.size() == 4 && in buildAtomicFloatingRMWInst()
847 "Wrong number of atomic floating-type builtin"); in buildAtomicFloatingRMWInst()
851 Register PtrReg = Call->Arguments[0]; in buildAtomicFloatingRMWInst()
852 MRI->setRegClass(PtrReg, &SPIRV::IDRegClass); in buildAtomicFloatingRMWInst()
854 Register ScopeReg = Call->Arguments[1]; in buildAtomicFloatingRMWInst()
855 MRI->setRegClass(ScopeReg, &SPIRV::IDRegClass); in buildAtomicFloatingRMWInst()
857 Register MemSemanticsReg = Call->Arguments[2]; in buildAtomicFloatingRMWInst()
858 MRI->setRegClass(MemSemanticsReg, &SPIRV::IDRegClass); in buildAtomicFloatingRMWInst()
860 Register ValueReg = Call->Arguments[3]; in buildAtomicFloatingRMWInst()
861 MRI->setRegClass(ValueReg, &SPIRV::IDRegClass); in buildAtomicFloatingRMWInst()
864 .addDef(Call->ReturnRegister) in buildAtomicFloatingRMWInst()
865 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in buildAtomicFloatingRMWInst()
879 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType); in buildAtomicFlagInst()
880 if (Call->isSpirvOp()) in buildAtomicFlagInst()
885 Register PtrRegister = Call->Arguments[0]; in buildAtomicFlagInst()
888 Call->Arguments.size() >= 2 ? Call->Arguments[1] : Register(); in buildAtomicFlagInst()
898 Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register(); in buildAtomicFlagInst()
904 MIB.addDef(Call->ReturnRegister).addUse(TypeReg); in buildAtomicFlagInst()
915 if (Call->isSpirvOp()) in buildBarrierInst()
919 unsigned MemFlags = getIConstVal(Call->Arguments[0], MRI); in buildBarrierInst()
933 static_cast<std::memory_order>(getIConstVal(Call->Arguments[1], MRI)); in buildBarrierInst()
941 MemSemanticsReg = Call->Arguments[0]; in buildBarrierInst()
942 MRI->setRegClass(MemSemanticsReg, &SPIRV::IDRegClass); in buildBarrierInst()
949 if (Call->Arguments.size() >= 2) { in buildBarrierInst()
951 ((Opcode != SPIRV::OpMemoryBarrier && Call->Arguments.size() == 2) || in buildBarrierInst()
952 (Opcode == SPIRV::OpMemoryBarrier && Call->Arguments.size() == 3)) && in buildBarrierInst()
954 Register ScopeArg = (Opcode == SPIRV::OpMemoryBarrier) ? Call->Arguments[2] in buildBarrierInst()
955 : Call->Arguments[1]; in buildBarrierInst()
964 ScopeReg = Call->Arguments[1]; in buildBarrierInst()
965 MRI->setRegClass(ScopeReg, &SPIRV::IDRegClass); in buildBarrierInst()
979 static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim) { in getNumComponentsForDim() argument
980 switch (dim) { in getNumComponentsForDim()
981 case SPIRV::Dim::DIM_1D: in getNumComponentsForDim()
982 case SPIRV::Dim::DIM_Buffer: in getNumComponentsForDim()
984 case SPIRV::Dim::DIM_2D: in getNumComponentsForDim()
985 case SPIRV::Dim::DIM_Cube: in getNumComponentsForDim()
986 case SPIRV::Dim::DIM_Rect: in getNumComponentsForDim()
988 case SPIRV::Dim::DIM_3D: in getNumComponentsForDim()
991 report_fatal_error("Cannot get num components for given Dim"); in getNumComponentsForDim()
997 assert(imgType->getOpcode() == SPIRV::OpTypeImage); in getNumSizeComponents()
998 auto dim = static_cast<SPIRV::Dim::Dim>(imgType->getOperand(2).getImm()); in getNumSizeComponents() local
999 unsigned numComps = getNumComponentsForDim(dim); in getNumSizeComponents()
1000 bool arrayed = imgType->getOperand(4).getImm() == 1; in getNumSizeComponents()
1004 //===----------------------------------------------------------------------===//
1006 //===----------------------------------------------------------------------===//
1012 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateExtInst()
1014 SPIRV::lookupExtendedBuiltin(Builtin->Name, Builtin->Set)->Number; in generateExtInst()
1019 .addDef(Call->ReturnRegister) in generateExtInst()
1020 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateExtInst()
1024 for (auto Argument : Call->Arguments) in generateExtInst()
1033 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateRelationalInst()
1035 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode; in generateRelationalInst()
1040 buildBoolRegister(MIRBuilder, Call->ReturnType, GR); in generateRelationalInst()
1045 .addUse(GR->getSPIRVTypeID(RelationType)); in generateRelationalInst()
1047 for (auto Argument : Call->Arguments) in generateRelationalInst()
1051 return buildSelectInst(MIRBuilder, Call->ReturnRegister, CompareRegister, in generateRelationalInst()
1052 Call->ReturnType, GR); in generateRelationalInst()
1058 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateGroupInst()
1060 SPIRV::lookupGroupBuiltin(Builtin->Name); in generateGroupInst()
1063 if (Call->isSpirvOp()) { in generateGroupInst()
1064 if (GroupBuiltin->NoGroupOperation) in generateGroupInst()
1065 return buildOpFromWrapper(MIRBuilder, GroupBuiltin->Opcode, Call, in generateGroupInst()
1066 GR->getSPIRVTypeID(Call->ReturnType)); in generateGroupInst()
1069 Register GroupOpReg = Call->Arguments[1]; in generateGroupInst()
1071 if (!MI || MI->getOpcode() != TargetOpcode::G_CONSTANT) in generateGroupInst()
1074 uint64_t GrpOp = MI->getOperand(1).getCImm()->getValue().getZExtValue(); in generateGroupInst()
1075 Register ScopeReg = Call->Arguments[0]; in generateGroupInst()
1076 if (!MRI->getRegClassOrNull(ScopeReg)) in generateGroupInst()
1077 MRI->setRegClass(ScopeReg, &SPIRV::IDRegClass); in generateGroupInst()
1078 auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode) in generateGroupInst()
1079 .addDef(Call->ReturnRegister) in generateGroupInst()
1080 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateGroupInst()
1083 for (unsigned i = 2; i < Call->Arguments.size(); ++i) { in generateGroupInst()
1084 Register ArgReg = Call->Arguments[i]; in generateGroupInst()
1085 if (!MRI->getRegClassOrNull(ArgReg)) in generateGroupInst()
1086 MRI->setRegClass(ArgReg, &SPIRV::IDRegClass); in generateGroupInst()
1093 if (GroupBuiltin->HasBoolArg) { in generateGroupInst()
1094 Register ConstRegister = Call->Arguments[0]; in generateGroupInst()
1097 // TODO: support non-constant bool values. in generateGroupInst()
1098 assert(ArgInstruction->getOpcode() == TargetOpcode::G_CONSTANT && in generateGroupInst()
1100 if (GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode() != in generateGroupInst()
1102 Arg0 = GR->buildConstantInt(getIConstVal(ConstRegister, MRI), MIRBuilder, in generateGroupInst()
1103 GR->getOrCreateSPIRVBoolType(MIRBuilder)); in generateGroupInst()
1106 Register GroupResultRegister = Call->ReturnRegister; in generateGroupInst()
1107 SPIRVType *GroupResultType = Call->ReturnType; in generateGroupInst()
1112 GroupBuiltin->IsElect || GroupBuiltin->IsAllOrAny || in generateGroupInst()
1113 GroupBuiltin->IsAllEqual || GroupBuiltin->IsLogical || in generateGroupInst()
1114 GroupBuiltin->IsInverseBallot || GroupBuiltin->IsBallotBitExtract; in generateGroupInst()
1118 buildBoolRegister(MIRBuilder, Call->ReturnType, GR); in generateGroupInst()
1120 auto Scope = Builtin->Name.starts_with("sub_group") ? SPIRV::Scope::Subgroup in generateGroupInst()
1125 auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode) in generateGroupInst()
1127 .addUse(GR->getSPIRVTypeID(GroupResultType)) in generateGroupInst()
1130 if (!GroupBuiltin->NoGroupOperation) in generateGroupInst()
1131 MIB.addImm(GroupBuiltin->GroupOperation); in generateGroupInst()
1132 if (Call->Arguments.size() > 0) { in generateGroupInst()
1133 MIB.addUse(Arg0.isValid() ? Arg0 : Call->Arguments[0]); in generateGroupInst()
1134 MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass); in generateGroupInst()
1135 for (unsigned i = 1; i < Call->Arguments.size(); i++) { in generateGroupInst()
1136 MIB.addUse(Call->Arguments[i]); in generateGroupInst()
1137 MRI->setRegClass(Call->Arguments[i], &SPIRV::IDRegClass); in generateGroupInst()
1143 buildSelectInst(MIRBuilder, Call->ReturnRegister, GroupResultRegister, in generateGroupInst()
1144 Call->ReturnType, GR); in generateGroupInst()
1151 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateIntelSubgroupsInst()
1154 if (!ST->canUseExtension(SPIRV::Extension::SPV_INTEL_subgroups)) { in generateIntelSubgroupsInst()
1155 std::string DiagMsg = std::string(Builtin->Name) + in generateIntelSubgroupsInst()
1156 ": the builtin requires the following SPIR-V " in generateIntelSubgroupsInst()
1161 SPIRV::lookupIntelSubgroupsBuiltin(Builtin->Name); in generateIntelSubgroupsInst()
1163 uint32_t OpCode = IntelSubgroups->Opcode; in generateIntelSubgroupsInst()
1164 if (Call->isSpirvOp()) { in generateIntelSubgroupsInst()
1168 IsSet ? GR->getSPIRVTypeID(Call->ReturnType) in generateIntelSubgroupsInst()
1173 if (IntelSubgroups->IsBlock) { in generateIntelSubgroupsInst()
1175 if (SPIRVType *Arg0Type = GR->getSPIRVTypeForVReg(Call->Arguments[0])) { in generateIntelSubgroupsInst()
1176 if (Arg0Type->getOpcode() == SPIRV::OpTypeImage) { in generateIntelSubgroupsInst()
1204 IntelSubgroups->IsWrite in generateIntelSubgroupsInst()
1207 .addDef(Call->ReturnRegister) in generateIntelSubgroupsInst()
1208 .addUse(GR->getSPIRVTypeID(Call->ReturnType)); in generateIntelSubgroupsInst()
1209 for (size_t i = 0; i < Call->Arguments.size(); ++i) { in generateIntelSubgroupsInst()
1210 MIB.addUse(Call->Arguments[i]); in generateIntelSubgroupsInst()
1211 MRI->setRegClass(Call->Arguments[i], &SPIRV::IDRegClass); in generateIntelSubgroupsInst()
1220 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateGroupUniformInst()
1223 if (!ST->canUseExtension( in generateGroupUniformInst()
1225 std::string DiagMsg = std::string(Builtin->Name) + in generateGroupUniformInst()
1226 ": the builtin requires the following SPIR-V " in generateGroupUniformInst()
1231 SPIRV::lookupGroupUniformBuiltin(Builtin->Name); in generateGroupUniformInst()
1234 Register GroupResultReg = Call->ReturnRegister; in generateGroupUniformInst()
1235 MRI->setRegClass(GroupResultReg, &SPIRV::IDRegClass); in generateGroupUniformInst()
1238 Register ScopeReg = Call->Arguments[0]; in generateGroupUniformInst()
1239 MRI->setRegClass(ScopeReg, &SPIRV::IDRegClass); in generateGroupUniformInst()
1242 Register ConstGroupOpReg = Call->Arguments[1]; in generateGroupUniformInst()
1244 if (!Const || Const->getOpcode() != TargetOpcode::G_CONSTANT) in generateGroupUniformInst()
1248 const MachineOperand &ConstOperand = Const->getOperand(1); in generateGroupUniformInst()
1255 Register ValueReg = Call->Arguments[2]; in generateGroupUniformInst()
1256 MRI->setRegClass(ValueReg, &SPIRV::IDRegClass); in generateGroupUniformInst()
1258 auto MIB = MIRBuilder.buildInstr(GroupUniform->Opcode) in generateGroupUniformInst()
1260 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateGroupUniformInst()
1262 addNumImm(ConstOperand.getCImm()->getValue(), MIB); in generateGroupUniformInst()
1271 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateKernelClockInst()
1274 if (!ST->canUseExtension(SPIRV::Extension::SPV_KHR_shader_clock)) { in generateKernelClockInst()
1275 std::string DiagMsg = std::string(Builtin->Name) + in generateKernelClockInst()
1276 ": the builtin requires the following SPIR-V " in generateKernelClockInst()
1282 Register ResultReg = Call->ReturnRegister; in generateKernelClockInst()
1283 MRI->setRegClass(ResultReg, &SPIRV::IDRegClass); in generateKernelClockInst()
1287 StringSwitch<SPIRV::Scope::Scope>(Builtin->Name) in generateKernelClockInst()
1295 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateKernelClockInst()
1302 // size_t get_global_id(uint dimindex). In SPIR-V, the builtins corresonding to
1332 Register IndexRegister = Call->Arguments[0]; in genWorkgroupQuery()
1333 const unsigned ResultWidth = Call->ReturnType->getOperand(1).getImm(); in genWorkgroupQuery()
1334 const unsigned PointerSize = GR->getPointerSize(); in genWorkgroupQuery()
1336 GR->getOrCreateSPIRVIntegerType(PointerSize, MIRBuilder); in genWorkgroupQuery()
1341 Register ToTruncate = Call->ReturnRegister; in genWorkgroupQuery()
1345 IndexInstruction->getOpcode() == TargetOpcode::G_CONSTANT; in genWorkgroupQuery()
1350 Register DefaultReg = Call->ReturnRegister; in genWorkgroupQuery()
1352 DefaultReg = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize)); in genWorkgroupQuery()
1353 MRI->setRegClass(DefaultReg, &SPIRV::IDRegClass); in genWorkgroupQuery()
1354 GR->assignSPIRVTypeToVReg(PointerSizeType, DefaultReg, in genWorkgroupQuery()
1359 GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType); in genWorkgroupQuery()
1363 GR->getOrCreateSPIRVVectorType(PointerSizeType, 3, MIRBuilder); in genWorkgroupQuery()
1368 Register Extracted = Call->ReturnRegister; in genWorkgroupQuery()
1370 Extracted = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize)); in genWorkgroupQuery()
1371 MRI->setRegClass(Extracted, &SPIRV::IDRegClass); in genWorkgroupQuery()
1372 GR->assignSPIRVTypeToVReg(PointerSizeType, Extracted, MIRBuilder.getMF()); in genWorkgroupQuery()
1385 auto IndexType = GR->getSPIRVTypeForVReg(IndexRegister); in genWorkgroupQuery()
1386 auto BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder); in genWorkgroupQuery()
1389 MRI->createGenericVirtualRegister(LLT::scalar(1)); in genWorkgroupQuery()
1390 MRI->setRegClass(CompareRegister, &SPIRV::IDRegClass); in genWorkgroupQuery()
1391 GR->assignSPIRVTypeToVReg(BoolType, CompareRegister, MIRBuilder.getMF()); in genWorkgroupQuery()
1395 GR->buildConstantInt(3, MIRBuilder, IndexType)); in genWorkgroupQuery()
1400 GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType); in genWorkgroupQuery()
1403 Register SelectionResult = Call->ReturnRegister; in genWorkgroupQuery()
1406 MRI->createGenericVirtualRegister(LLT::scalar(PointerSize)); in genWorkgroupQuery()
1407 MRI->setRegClass(SelectionResult, &SPIRV::IDRegClass); in genWorkgroupQuery()
1408 GR->assignSPIRVTypeToVReg(PointerSizeType, SelectionResult, in genWorkgroupQuery()
1421 MIRBuilder.buildZExtOrTrunc(Call->ReturnRegister, ToTruncate); in genWorkgroupQuery()
1429 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateBuiltinVar()
1431 SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value; in generateBuiltinVar()
1437 unsigned BitWidth = GR->getScalarOrVectorBitWidth(Call->ReturnType); in generateBuiltinVar()
1439 if (Call->ReturnType->getOpcode() == SPIRV::OpTypeVector) in generateBuiltinVar()
1441 LLT::fixed_vector(Call->ReturnType->getOperand(2).getImm(), BitWidth); in generateBuiltinVar()
1445 return buildBuiltinVariableLoad(MIRBuilder, Call->ReturnType, GR, Value, in generateBuiltinVar()
1446 LLType, Call->ReturnRegister); in generateBuiltinVar()
1453 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateAtomicInst()
1455 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode; in generateAtomicInst()
1481 if (Call->isSpirvOp()) in generateAtomicInst()
1483 GR->getSPIRVTypeID(Call->ReturnType)); in generateAtomicInst()
1492 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateAtomicFloatingInst()
1493 unsigned Opcode = SPIRV::lookupAtomicFloatingBuiltin(Builtin->Name)->Opcode; in generateAtomicFloatingInst()
1509 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateBarrierInst()
1511 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode; in generateBarrierInst()
1519 .addDef(Call->ReturnRegister) in generateCastToPtrInst()
1520 .addUse(Call->Arguments[0]); in generateCastToPtrInst()
1527 if (Call->isSpirvOp()) in generateDotOrFMulInst()
1529 GR->getSPIRVTypeID(Call->ReturnType)); in generateDotOrFMulInst()
1530 unsigned Opcode = GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode(); in generateDotOrFMulInst()
1534 .addDef(Call->ReturnRegister) in generateDotOrFMulInst()
1535 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateDotOrFMulInst()
1536 .addUse(Call->Arguments[0]) in generateDotOrFMulInst()
1537 .addUse(Call->Arguments[1]); in generateDotOrFMulInst()
1544 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateWaveInst()
1546 SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value; in generateWaveInst()
1549 assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt); in generateWaveInst()
1550 LLT LLType = LLT::scalar(GR->getScalarOrVectorBitWidth(Call->ReturnType)); in generateWaveInst()
1553 MIRBuilder, Call->ReturnType, GR, Value, LLType, Call->ReturnRegister, in generateWaveInst()
1562 SPIRV::lookupGetBuiltin(Call->Builtin->Name, Call->Builtin->Set)->Value; in generateGetQueryInst()
1573 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateImageSizeQueryInst()
1575 SPIRV::lookupImageQueryBuiltin(Builtin->Name, Builtin->Set)->Component; in generateImageSizeQueryInst()
1579 SPIRVType *RetTy = Call->ReturnType; in generateImageSizeQueryInst()
1580 unsigned NumExpectedRetComponents = RetTy->getOpcode() == SPIRV::OpTypeVector in generateImageSizeQueryInst()
1581 ? RetTy->getOperand(2).getImm() in generateImageSizeQueryInst()
1584 SPIRVType *ImgType = GR->getSPIRVTypeForVReg(Call->Arguments[0]); in generateImageSizeQueryInst()
1586 Register QueryResult = Call->ReturnRegister; in generateImageSizeQueryInst()
1587 SPIRVType *QueryResultType = Call->ReturnType; in generateImageSizeQueryInst()
1589 QueryResult = MIRBuilder.getMRI()->createGenericVirtualRegister( in generateImageSizeQueryInst()
1591 MIRBuilder.getMRI()->setRegClass(QueryResult, &SPIRV::IDRegClass); in generateImageSizeQueryInst()
1592 SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder); in generateImageSizeQueryInst()
1593 QueryResultType = GR->getOrCreateSPIRVVectorType( in generateImageSizeQueryInst()
1595 GR->assignSPIRVTypeToVReg(QueryResultType, QueryResult, MIRBuilder.getMF()); in generateImageSizeQueryInst()
1597 bool IsDimBuf = ImgType->getOperand(2).getImm() == SPIRV::Dim::DIM_Buffer; in generateImageSizeQueryInst()
1600 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass); in generateImageSizeQueryInst()
1603 .addUse(GR->getSPIRVTypeID(QueryResultType)) in generateImageSizeQueryInst()
1604 .addUse(Call->Arguments[0]); in generateImageSizeQueryInst()
1612 Component == 3 ? NumActualRetComponents - 1 : Component; in generateImageSizeQueryInst()
1615 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType); in generateImageSizeQueryInst()
1617 if (QueryResultType->getOpcode() == SPIRV::OpTypeVector) { in generateImageSizeQueryInst()
1618 Register NewTypeReg = QueryResultType->getOperand(1).getReg(); in generateImageSizeQueryInst()
1620 (NewType = GR->getSPIRVTypeForVReg(NewTypeReg)) != nullptr) in generateImageSizeQueryInst()
1624 .addDef(Call->ReturnRegister) in generateImageSizeQueryInst()
1629 insertAssignInstr(Call->ReturnRegister, nullptr, NewType, GR, MIRBuilder, in generateImageSizeQueryInst()
1634 .addDef(Call->ReturnRegister) in generateImageSizeQueryInst()
1635 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateImageSizeQueryInst()
1647 assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt && in generateImageMiscQueryInst()
1651 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateImageMiscQueryInst()
1653 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode; in generateImageMiscQueryInst()
1655 Register Image = Call->Arguments[0]; in generateImageMiscQueryInst()
1656 MIRBuilder.getMRI()->setRegClass(Image, &SPIRV::IDRegClass); in generateImageMiscQueryInst()
1657 SPIRV::Dim::Dim ImageDimensionality = static_cast<SPIRV::Dim::Dim>( in generateImageMiscQueryInst()
1658 GR->getSPIRVTypeForVReg(Image)->getOperand(2).getImm()); in generateImageMiscQueryInst()
1663 assert(ImageDimensionality == SPIRV::Dim::DIM_2D && in generateImageMiscQueryInst()
1667 assert((ImageDimensionality == SPIRV::Dim::DIM_1D || in generateImageMiscQueryInst()
1668 ImageDimensionality == SPIRV::Dim::DIM_2D || in generateImageMiscQueryInst()
1669 ImageDimensionality == SPIRV::Dim::DIM_3D || in generateImageMiscQueryInst()
1670 ImageDimensionality == SPIRV::Dim::DIM_Cube) && in generateImageMiscQueryInst()
1676 .addDef(Call->ReturnRegister) in generateImageMiscQueryInst()
1677 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateImageMiscQueryInst()
1697 report_fatal_error("Unknown CL address mode"); in getSamplerAddressingModeFromBitmask()
1718 Register Image = Call->Arguments[0]; in generateReadImageInst()
1720 MRI->setRegClass(Image, &SPIRV::IDRegClass); in generateReadImageInst()
1721 MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass); in generateReadImageInst()
1725 MRI->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass); in generateReadImageInst()
1727 Register Sampler = Call->Arguments[1]; in generateReadImageInst()
1729 if (!GR->isScalarOfType(Sampler, SPIRV::OpTypeSampler) && in generateReadImageInst()
1730 getDefInstrMaybeConstant(Sampler, MRI)->getOperand(1).isCImm()) { in generateReadImageInst()
1732 Sampler = GR->buildConstantSampler( in generateReadImageInst()
1736 GR->getSPIRVTypeForVReg(Sampler)); in generateReadImageInst()
1738 SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image); in generateReadImageInst()
1740 GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder); in generateReadImageInst()
1741 Register SampledImage = MRI->createVirtualRegister(&SPIRV::IDRegClass); in generateReadImageInst()
1745 .addUse(GR->getSPIRVTypeID(SampledImageType)) in generateReadImageInst()
1749 Register Lod = GR->buildConstantFP(APFloat::getZero(APFloat::IEEEsingle()), in generateReadImageInst()
1751 SPIRVType *TempType = Call->ReturnType; in generateReadImageInst()
1753 if (TempType->getOpcode() != SPIRV::OpTypeVector) { in generateReadImageInst()
1755 GR->getOrCreateSPIRVVectorType(Call->ReturnType, 4, MIRBuilder); in generateReadImageInst()
1758 LLT LLType = LLT::scalar(GR->getScalarOrVectorBitWidth(TempType)); in generateReadImageInst()
1759 Register TempRegister = MRI->createGenericVirtualRegister(LLType); in generateReadImageInst()
1760 MRI->setRegClass(TempRegister, &SPIRV::IDRegClass); in generateReadImageInst()
1761 GR->assignSPIRVTypeToVReg(TempType, TempRegister, MIRBuilder.getMF()); in generateReadImageInst()
1764 .addDef(NeedsExtraction ? TempRegister : Call->ReturnRegister) in generateReadImageInst()
1765 .addUse(GR->getSPIRVTypeID(TempType)) in generateReadImageInst()
1767 .addUse(Call->Arguments[2]) // Coordinate. in generateReadImageInst()
1773 .addDef(Call->ReturnRegister) in generateReadImageInst()
1774 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateReadImageInst()
1779 .addDef(Call->ReturnRegister) in generateReadImageInst()
1780 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateReadImageInst()
1782 .addUse(Call->Arguments[1]) // Coordinate. in generateReadImageInst()
1784 .addUse(Call->Arguments[2]); in generateReadImageInst()
1787 .addDef(Call->ReturnRegister) in generateReadImageInst()
1788 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateReadImageInst()
1790 .addUse(Call->Arguments[1]); // Coordinate. in generateReadImageInst()
1798 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass); in generateWriteImageInst()
1799 MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass); in generateWriteImageInst()
1800 MIRBuilder.getMRI()->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass); in generateWriteImageInst()
1802 .addUse(Call->Arguments[0]) // Image. in generateWriteImageInst()
1803 .addUse(Call->Arguments[1]) // Coordinate. in generateWriteImageInst()
1804 .addUse(Call->Arguments[2]); // Texel. in generateWriteImageInst()
1813 if (Call->Builtin->Name.contains_insensitive( in generateSampleImageInst()
1816 uint64_t Bitmask = getIConstVal(Call->Arguments[0], MRI); in generateSampleImageInst()
1817 Register Sampler = GR->buildConstantSampler( in generateSampleImageInst()
1818 Call->ReturnRegister, getSamplerAddressingModeFromBitmask(Bitmask), in generateSampleImageInst()
1820 getSamplerFilterModeFromBitmask(Bitmask), MIRBuilder, Call->ReturnType); in generateSampleImageInst()
1822 } else if (Call->Builtin->Name.contains_insensitive("__spirv_SampledImage")) { in generateSampleImageInst()
1824 Register Image = Call->Arguments[0]; in generateSampleImageInst()
1825 SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image); in generateSampleImageInst()
1827 GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder); in generateSampleImageInst()
1829 Call->ReturnRegister.isValid() in generateSampleImageInst()
1830 ? Call->ReturnRegister in generateSampleImageInst()
1831 : MRI->createVirtualRegister(&SPIRV::IDRegClass); in generateSampleImageInst()
1834 .addUse(GR->getSPIRVTypeID(SampledImageType)) in generateSampleImageInst()
1836 .addUse(Call->Arguments[1]); // Sampler. in generateSampleImageInst()
1838 } else if (Call->Builtin->Name.contains_insensitive( in generateSampleImageInst()
1847 Call->ReturnType in generateSampleImageInst()
1848 ? Call->ReturnType in generateSampleImageInst()
1849 : GR->getOrCreateSPIRVTypeByName(ReturnType, MIRBuilder); in generateSampleImageInst()
1855 MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass); in generateSampleImageInst()
1856 MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass); in generateSampleImageInst()
1857 MRI->setRegClass(Call->Arguments[3], &SPIRV::IDRegClass); in generateSampleImageInst()
1860 .addDef(Call->ReturnRegister) in generateSampleImageInst()
1861 .addUse(GR->getSPIRVTypeID(Type)) in generateSampleImageInst()
1862 .addUse(Call->Arguments[0]) // Image. in generateSampleImageInst()
1863 .addUse(Call->Arguments[1]) // Coordinate. in generateSampleImageInst()
1865 .addUse(Call->Arguments[3]); in generateSampleImageInst()
1873 MIRBuilder.buildSelect(Call->ReturnRegister, Call->Arguments[0], in generateSelectInst()
1874 Call->Arguments[1], Call->Arguments[2]); in generateSelectInst()
1882 GR->getSPIRVTypeID(Call->ReturnType)); in generateConstructInst()
1888 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateCoopMatrInst()
1890 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode; in generateCoopMatrInst()
1892 unsigned ArgSz = Call->Arguments.size(); in generateCoopMatrInst()
1901 ImmArgs.push_back(getConstFromIntrinsic(Call->Arguments[LiteralIdx], MRI)); in generateCoopMatrInst()
1902 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType); in generateCoopMatrInst()
1904 SPIRVType *CoopMatrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]); in generateCoopMatrInst()
1908 .addDef(Call->ReturnRegister) in generateCoopMatrInst()
1910 .addUse(CoopMatrType->getOperand(0).getReg()); in generateCoopMatrInst()
1921 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateSpecConstantInst()
1923 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode; in generateSpecConstantInst()
1930 static_cast<unsigned>(getIConstVal(Call->Arguments[0], MRI)); in generateSpecConstantInst()
1931 buildOpDecorate(Call->ReturnRegister, MIRBuilder, SPIRV::Decoration::SpecId, in generateSpecConstantInst()
1934 Register ConstRegister = Call->Arguments[1]; in generateSpecConstantInst()
1937 (Const->getOpcode() == TargetOpcode::G_CONSTANT || in generateSpecConstantInst()
1938 Const->getOpcode() == TargetOpcode::G_FCONSTANT) && in generateSpecConstantInst()
1939 "Argument should be either an int or floating-point constant"); in generateSpecConstantInst()
1941 const MachineOperand &ConstOperand = Const->getOperand(1); in generateSpecConstantInst()
1942 if (Call->ReturnType->getOpcode() == SPIRV::OpTypeBool) { in generateSpecConstantInst()
1944 Opcode = ConstOperand.getCImm()->getValue().getZExtValue() in generateSpecConstantInst()
1949 .addDef(Call->ReturnRegister) in generateSpecConstantInst()
1950 .addUse(GR->getSPIRVTypeID(Call->ReturnType)); in generateSpecConstantInst()
1952 if (Call->ReturnType->getOpcode() != SPIRV::OpTypeBool) { in generateSpecConstantInst()
1953 if (Const->getOpcode() == TargetOpcode::G_CONSTANT) in generateSpecConstantInst()
1954 addNumImm(ConstOperand.getCImm()->getValue(), MIB); in generateSpecConstantInst()
1956 addNumImm(ConstOperand.getFPImm()->getValueAPF().bitcastToAPInt(), MIB); in generateSpecConstantInst()
1962 .addDef(Call->ReturnRegister) in generateSpecConstantInst()
1963 .addUse(GR->getSPIRVTypeID(Call->ReturnType)); in generateSpecConstantInst()
1964 for (unsigned i = 0; i < Call->Arguments.size(); i++) in generateSpecConstantInst()
1965 MIB.addUse(Call->Arguments[i]); in generateSpecConstantInst()
1977 MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass); in buildNDRange()
1978 SPIRVType *PtrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]); in buildNDRange()
1979 assert(PtrType->getOpcode() == SPIRV::OpTypePointer && in buildNDRange()
1980 PtrType->getOperand(2).isReg()); in buildNDRange()
1981 Register TypeReg = PtrType->getOperand(2).getReg(); in buildNDRange()
1982 SPIRVType *StructType = GR->getSPIRVTypeForVReg(TypeReg); in buildNDRange()
1984 Register TmpReg = MRI->createVirtualRegister(&SPIRV::IDRegClass); in buildNDRange()
1985 GR->assignSPIRVTypeToVReg(StructType, TmpReg, MF); in buildNDRange()
1988 unsigned NumArgs = Call->Arguments.size(); in buildNDRange()
1990 Register GlobalWorkSize = Call->Arguments[NumArgs < 4 ? 1 : 2]; in buildNDRange()
1991 MRI->setRegClass(GlobalWorkSize, &SPIRV::IDRegClass); in buildNDRange()
1993 NumArgs == 2 ? Register(0) : Call->Arguments[NumArgs < 4 ? 2 : 3]; in buildNDRange()
1995 MRI->setRegClass(LocalWorkSize, &SPIRV::IDRegClass); in buildNDRange()
1996 Register GlobalWorkOffset = NumArgs <= 3 ? Register(0) : Call->Arguments[1]; in buildNDRange()
1998 MRI->setRegClass(GlobalWorkOffset, &SPIRV::IDRegClass); in buildNDRange()
2001 SPIRVType *SpvTy = GR->getSPIRVTypeForVReg(GlobalWorkSize); in buildNDRange()
2002 if (SpvTy->getOpcode() == SPIRV::OpTypePointer) { in buildNDRange()
2003 MachineInstr *DefInstr = MRI->getUniqueVRegDef(GlobalWorkSize); in buildNDRange()
2005 DefInstr->getOperand(3).isReg()); in buildNDRange()
2006 Register GWSPtr = DefInstr->getOperand(3).getReg(); in buildNDRange()
2007 if (!MRI->getRegClassOrNull(GWSPtr)) in buildNDRange()
2008 MRI->setRegClass(GWSPtr, &SPIRV::IDRegClass); in buildNDRange()
2010 unsigned Size = Call->Builtin->Name == "ndrange_3D" ? 3 : 2; in buildNDRange()
2011 unsigned BitWidth = GR->getPointerSize() == 64 ? 64 : 32; in buildNDRange()
2014 SPIRVType *SpvFieldTy = GR->getOrCreateSPIRVType(FieldTy, MIRBuilder); in buildNDRange()
2015 GlobalWorkSize = MRI->createVirtualRegister(&SPIRV::IDRegClass); in buildNDRange()
2016 GR->assignSPIRVTypeToVReg(SpvFieldTy, GlobalWorkSize, MF); in buildNDRange()
2019 .addUse(GR->getSPIRVTypeID(SpvFieldTy)) in buildNDRange()
2023 Const = GR->getOrCreateConstIntArray(0, Size, *MIRBuilder.getInsertPt(), in buildNDRange()
2026 Const = GR->buildConstantInt(0, MIRBuilder, SpvTy); in buildNDRange()
2041 .addUse(Call->Arguments[0]) in buildNDRange()
2058 return GR->getOrCreateSPIRVType(PtrType, MIRBuilder); in getOrCreateSPIRVDeviceEventPointer()
2066 bool IsSpirvOp = Call->isSpirvOp(); in buildEnqueueKernel()
2067 bool HasEvents = Call->Builtin->Name.contains("events") || IsSpirvOp; in buildEnqueueKernel()
2068 const SPIRVType *Int32Ty = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder); in buildEnqueueKernel()
2074 if (Call->Builtin->Name.contains("_varargs") || IsSpirvOp) { in buildEnqueueKernel()
2076 Register GepReg = Call->Arguments[LocalSizeArrayIdx]; in buildEnqueueKernel()
2077 MachineInstr *GepMI = MRI->getUniqueVRegDef(GepReg); in buildEnqueueKernel()
2079 GepMI->getOperand(3).isReg()); in buildEnqueueKernel()
2080 Register ArrayReg = GepMI->getOperand(3).getReg(); in buildEnqueueKernel()
2081 MachineInstr *ArrayMI = MRI->getUniqueVRegDef(ArrayReg); in buildEnqueueKernel()
2085 cast<ArrayType>(LocalSizeTy)->getNumElements(); in buildEnqueueKernel()
2087 const LLT LLType = LLT::pointer(SC, GR->getPointerSize()); in buildEnqueueKernel()
2088 const SPIRVType *PointerSizeTy = GR->getOrCreateSPIRVPointerType( in buildEnqueueKernel()
2091 Register Reg = MRI->createVirtualRegister(&SPIRV::IDRegClass); in buildEnqueueKernel()
2092 MRI->setType(Reg, LLType); in buildEnqueueKernel()
2093 GR->assignSPIRVTypeToVReg(PointerSizeTy, Reg, MIRBuilder.getMF()); in buildEnqueueKernel()
2097 .addImm(GepMI->getOperand(2).getImm()) // In bound. in buildEnqueueKernel()
2098 .addUse(ArrayMI->getOperand(0).getReg()) // Alloca. in buildEnqueueKernel()
2107 .addDef(Call->ReturnRegister) in buildEnqueueKernel()
2108 .addUse(GR->getSPIRVTypeID(Int32Ty)); in buildEnqueueKernel()
2113 MIB.addUse(Call->Arguments[i]); in buildEnqueueKernel()
2118 Register NullPtr = GR->getOrCreateConstNullPtr( in buildEnqueueKernel()
2124 MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI); in buildEnqueueKernel()
2125 assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE); in buildEnqueueKernel()
2127 MIB.addGlobalAddress(BlockMI->getOperand(1).getGlobal()); in buildEnqueueKernel()
2129 Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1]; in buildEnqueueKernel()
2150 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateEnqueueInst()
2152 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode; in generateEnqueueInst()
2157 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass); in generateEnqueueInst()
2158 return MIRBuilder.buildInstr(Opcode).addUse(Call->Arguments[0]); in generateEnqueueInst()
2162 .addDef(Call->ReturnRegister) in generateEnqueueInst()
2163 .addUse(GR->getSPIRVTypeID(Call->ReturnType)); in generateEnqueueInst()
2165 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass); in generateEnqueueInst()
2167 .addDef(Call->ReturnRegister) in generateEnqueueInst()
2168 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateEnqueueInst()
2169 .addUse(Call->Arguments[0]); in generateEnqueueInst()
2171 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass); in generateEnqueueInst()
2172 MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass); in generateEnqueueInst()
2174 .addUse(Call->Arguments[0]) in generateEnqueueInst()
2175 .addUse(Call->Arguments[1]); in generateEnqueueInst()
2177 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass); in generateEnqueueInst()
2178 MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass); in generateEnqueueInst()
2179 MIRBuilder.getMRI()->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass); in generateEnqueueInst()
2181 .addUse(Call->Arguments[0]) in generateEnqueueInst()
2182 .addUse(Call->Arguments[1]) in generateEnqueueInst()
2183 .addUse(Call->Arguments[2]); in generateEnqueueInst()
2197 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateAsyncCopy()
2199 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode; in generateAsyncCopy()
2202 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType); in generateAsyncCopy()
2203 if (Call->isSpirvOp()) in generateAsyncCopy()
2212 Call->ReturnType->getOpcode() == SPIRV::OpTypeEvent in generateAsyncCopy()
2214 : GR->getOrCreateSPIRVTypeByName("spirv.Event", MIRBuilder); in generateAsyncCopy()
2215 Register TypeReg = GR->getSPIRVTypeID(NewType ? NewType : Call->ReturnType); in generateAsyncCopy()
2216 unsigned NumArgs = Call->Arguments.size(); in generateAsyncCopy()
2217 Register EventReg = Call->Arguments[NumArgs - 1]; in generateAsyncCopy()
2219 .addDef(Call->ReturnRegister) in generateAsyncCopy()
2222 .addUse(Call->Arguments[0]) in generateAsyncCopy()
2223 .addUse(Call->Arguments[1]) in generateAsyncCopy()
2224 .addUse(Call->Arguments[2]) in generateAsyncCopy()
2225 .addUse(Call->Arguments.size() > 4 in generateAsyncCopy()
2226 ? Call->Arguments[3] in generateAsyncCopy()
2230 insertAssignInstr(Call->ReturnRegister, nullptr, NewType, GR, MIRBuilder, in generateAsyncCopy()
2237 .addUse(Call->Arguments[0]) in generateAsyncCopy()
2238 .addUse(Call->Arguments[1]); in generateAsyncCopy()
2250 SPIRV::lookupConvertBuiltin(Call->Builtin->Name, Call->Builtin->Set); in generateConvertInst()
2252 if (!Builtin && Call->isSpirvOp()) { in generateConvertInst()
2253 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateConvertInst()
2255 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode; in generateConvertInst()
2257 GR->getSPIRVTypeID(Call->ReturnType)); in generateConvertInst()
2260 if (Builtin->IsSaturated) in generateConvertInst()
2261 buildOpDecorate(Call->ReturnRegister, MIRBuilder, in generateConvertInst()
2263 if (Builtin->IsRounded) in generateConvertInst()
2264 buildOpDecorate(Call->ReturnRegister, MIRBuilder, in generateConvertInst()
2266 {(unsigned)Builtin->RoundingMode}); in generateConvertInst()
2271 if (GR->isScalarOrVectorOfType(Call->Arguments[0], SPIRV::OpTypeInt)) { in generateConvertInst()
2272 // Int -> ... in generateConvertInst()
2273 if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) { in generateConvertInst()
2274 // Int -> Int in generateConvertInst()
2275 if (Builtin->IsSaturated) in generateConvertInst()
2276 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpSatConvertUToS in generateConvertInst()
2279 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpUConvert in generateConvertInst()
2281 } else if (GR->isScalarOrVectorOfType(Call->ReturnRegister, in generateConvertInst()
2283 // Int -> Float in generateConvertInst()
2284 if (Builtin->IsBfloat16) { in generateConvertInst()
2287 if (!ST->canUseExtension( in generateConvertInst()
2291 GR->getScalarOrVectorComponentCount(Call->Arguments[0]) == in generateConvertInst()
2292 GR->getScalarOrVectorComponentCount(Call->ReturnRegister); in generateConvertInst()
2300 } else if (GR->isScalarOrVectorOfType(Call->Arguments[0], in generateConvertInst()
2302 // Float -> ... in generateConvertInst()
2303 if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) { in generateConvertInst()
2304 // Float -> Int in generateConvertInst()
2305 if (Builtin->IsBfloat16) { in generateConvertInst()
2308 if (!ST->canUseExtension( in generateConvertInst()
2312 GR->getScalarOrVectorComponentCount(Call->Arguments[0]) == in generateConvertInst()
2313 GR->getScalarOrVectorComponentCount(Call->ReturnRegister); in generateConvertInst()
2316 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpConvertFToS in generateConvertInst()
2319 } else if (GR->isScalarOrVectorOfType(Call->ReturnRegister, in generateConvertInst()
2321 // Float -> Float in generateConvertInst()
2327 std::string DiagMsg = std::string(Builtin->Name) + in generateConvertInst()
2328 ": the builtin requires the following SPIR-V " in generateConvertInst()
2335 std::string(Builtin->Name) + in generateConvertInst()
2343 .addDef(Call->ReturnRegister) in generateConvertInst()
2344 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateConvertInst()
2345 .addUse(Call->Arguments[0]); in generateConvertInst()
2354 SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name, in generateVectorLoadStoreInst()
2355 Call->Builtin->Set); in generateVectorLoadStoreInst()
2359 .addDef(Call->ReturnRegister) in generateVectorLoadStoreInst()
2360 .addUse(GR->getSPIRVTypeID(Call->ReturnType)) in generateVectorLoadStoreInst()
2362 .addImm(Builtin->Number); in generateVectorLoadStoreInst()
2363 for (auto Argument : Call->Arguments) in generateVectorLoadStoreInst()
2365 if (Builtin->Name.contains("load") && Builtin->ElementCount > 1) in generateVectorLoadStoreInst()
2366 MIB.addImm(Builtin->ElementCount); in generateVectorLoadStoreInst()
2368 // Rounding mode should be passed as a last argument in the MI for builtins in generateVectorLoadStoreInst()
2370 if (Builtin->IsRounded) in generateVectorLoadStoreInst()
2371 MIB.addImm(static_cast<uint32_t>(Builtin->RoundingMode)); in generateVectorLoadStoreInst()
2379 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin; in generateLoadStoreInst()
2381 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode; in generateLoadStoreInst()
2386 MIB.addDef(Call->ReturnRegister); in generateLoadStoreInst()
2387 MIB.addUse(GR->getSPIRVTypeID(Call->ReturnType)); in generateLoadStoreInst()
2390 MIB.addUse(Call->Arguments[0]); in generateLoadStoreInst()
2392 MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass); in generateLoadStoreInst()
2395 MIB.addUse(Call->Arguments[1]); in generateLoadStoreInst()
2396 MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass); in generateLoadStoreInst()
2399 unsigned NumArgs = Call->Arguments.size(); in generateLoadStoreInst()
2401 MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 1 : 2], MRI)); in generateLoadStoreInst()
2402 MRI->setRegClass(Call->Arguments[IsLoad ? 1 : 2], &SPIRV::IDRegClass); in generateLoadStoreInst()
2405 MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 2 : 3], MRI)); in generateLoadStoreInst()
2406 MRI->setRegClass(Call->Arguments[IsLoad ? 2 : 3], &SPIRV::IDRegClass); in generateLoadStoreInst()
2414 // tuple value <-1, 0, 0> if the builtin function is not found.
2415 // Not all builtin functions are supported, only those with a ready-to-use op
2427 return std::make_tuple(-1, 0, 0); in mapBuiltinToOpcode()
2429 switch (Call->Builtin->Group) { in mapBuiltinToOpcode()
2441 SPIRV::lookupNativeBuiltin(Call->Builtin->Name, Call->Builtin->Set)) in mapBuiltinToOpcode()
2442 return std::make_tuple(Call->Builtin->Group, R->Opcode, 0); in mapBuiltinToOpcode()
2445 if (const auto *R = SPIRV::lookupExtendedBuiltin(Call->Builtin->Name, in mapBuiltinToOpcode()
2446 Call->Builtin->Set)) in mapBuiltinToOpcode()
2447 return std::make_tuple(Call->Builtin->Group, 0, R->Number); in mapBuiltinToOpcode()
2450 if (const auto *R = SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name, in mapBuiltinToOpcode()
2451 Call->Builtin->Set)) in mapBuiltinToOpcode()
2452 return std::make_tuple(SPIRV::Extended, 0, R->Number); in mapBuiltinToOpcode()
2455 if (const auto *R = SPIRV::lookupGroupBuiltin(Call->Builtin->Name)) in mapBuiltinToOpcode()
2456 return std::make_tuple(Call->Builtin->Group, R->Opcode, 0); in mapBuiltinToOpcode()
2459 if (const auto *R = SPIRV::lookupAtomicFloatingBuiltin(Call->Builtin->Name)) in mapBuiltinToOpcode()
2460 return std::make_tuple(Call->Builtin->Group, R->Opcode, 0); in mapBuiltinToOpcode()
2463 if (const auto *R = SPIRV::lookupIntelSubgroupsBuiltin(Call->Builtin->Name)) in mapBuiltinToOpcode()
2464 return std::make_tuple(Call->Builtin->Group, R->Opcode, 0); in mapBuiltinToOpcode()
2467 if (const auto *R = SPIRV::lookupGroupUniformBuiltin(Call->Builtin->Name)) in mapBuiltinToOpcode()
2468 return std::make_tuple(Call->Builtin->Group, R->Opcode, 0); in mapBuiltinToOpcode()
2471 return std::make_tuple(Call->Builtin->Group, SPIRV::OpImageWrite, 0); in mapBuiltinToOpcode()
2473 return std::make_tuple(Call->Builtin->Group, TargetOpcode::G_SELECT, 0); in mapBuiltinToOpcode()
2475 return std::make_tuple(Call->Builtin->Group, SPIRV::OpCompositeConstruct, in mapBuiltinToOpcode()
2478 return std::make_tuple(Call->Builtin->Group, SPIRV::OpReadClockKHR, 0); in mapBuiltinToOpcode()
2480 return std::make_tuple(-1, 0, 0); in mapBuiltinToOpcode()
2482 return std::make_tuple(-1, 0, 0); in mapBuiltinToOpcode()
2493 // SPIR-V type and return register. in lowerBuiltin()
2496 if (OrigRetTy && !OrigRetTy->isVoidTy()) { in lowerBuiltin()
2497 ReturnType = GR->assignTypeToVReg(OrigRetTy, OrigRet, MIRBuilder); in lowerBuiltin()
2498 if (!MIRBuilder.getMRI()->getRegClassOrNull(ReturnRegister)) in lowerBuiltin()
2499 MIRBuilder.getMRI()->setRegClass(ReturnRegister, &SPIRV::IDRegClass); in lowerBuiltin()
2500 } else if (OrigRetTy && OrigRetTy->isVoidTy()) { in lowerBuiltin()
2501 ReturnRegister = MIRBuilder.getMRI()->createVirtualRegister(&IDRegClass); in lowerBuiltin()
2502 MIRBuilder.getMRI()->setType(ReturnRegister, LLT::scalar(32)); in lowerBuiltin()
2503 ReturnType = GR->assignTypeToVReg(OrigRetTy, ReturnRegister, MIRBuilder); in lowerBuiltin()
2516 assert(Args.size() >= Call->Builtin->MinNumArgs && in lowerBuiltin()
2518 if (Call->Builtin->MaxNumArgs && Args.size() > Call->Builtin->MaxNumArgs) in lowerBuiltin()
2522 switch (Call->Builtin->Group) { in lowerBuiltin()
2588 BuiltinArgs.split(BuiltinArgsTypeStrs, ',', -1, false); in parseBuiltinCallArgumentBaseType()
2631 BaseType->isVoidTy() ? Type::getInt8Ty(Ctx) : BaseType, VecElts, false); in parseBuiltinCallArgumentBaseType()
2655 //===----------------------------------------------------------------------===//
2657 //===----------------------------------------------------------------------===//
2671 //===----------------------------------------------------------------------===//
2673 //===----------------------------------------------------------------------===//
2679 unsigned Opcode = TypeRecord->Opcode; in getNonParameterizedType()
2681 return GR->getOrCreateOpTypeByOpcode(ExtensionType, MIRBuilder, Opcode); in getNonParameterizedType()
2687 return GR->getOrCreateOpTypeSampler(MIRBuilder); in getSamplerType()
2693 assert(ExtensionType->getNumIntParameters() == 1 && in getPipeType()
2694 "Invalid number of parameters for SPIR-V pipe builtin!"); in getPipeType()
2696 return GR->getOrCreateOpTypePipe(MIRBuilder, in getPipeType()
2698 ExtensionType->getIntParameter(0))); in getPipeType()
2704 assert(ExtensionType->getNumIntParameters() == 4 && in getCoopMatrType()
2705 "Invalid number of parameters for SPIR-V coop matrices builtin!"); in getCoopMatrType()
2706 assert(ExtensionType->getNumTypeParameters() == 1 && in getCoopMatrType()
2707 "SPIR-V coop matrices builtin type must have a type parameter!"); in getCoopMatrType()
2709 GR->getOrCreateSPIRVType(ExtensionType->getTypeParameter(0), MIRBuilder); in getCoopMatrType()
2711 return GR->getOrCreateOpTypeCoopMatr( in getCoopMatrType()
2712 MIRBuilder, ExtensionType, ElemType, ExtensionType->getIntParameter(0), in getCoopMatrType()
2713 ExtensionType->getIntParameter(1), ExtensionType->getIntParameter(2), in getCoopMatrType()
2714 ExtensionType->getIntParameter(3)); in getCoopMatrType()
2721 assert(ExtensionType->getNumTypeParameters() == 1 && in getImageType()
2722 "SPIR-V image builtin type must have sampled type parameter!"); in getImageType()
2724 GR->getOrCreateSPIRVType(ExtensionType->getTypeParameter(0), MIRBuilder); in getImageType()
2725 assert(ExtensionType->getNumIntParameters() == 7 && in getImageType()
2726 "Invalid number of parameters for SPIR-V image builtin!"); in getImageType()
2728 return GR->getOrCreateOpTypeImage( in getImageType()
2730 SPIRV::Dim::Dim(ExtensionType->getIntParameter(0)), in getImageType()
2731 ExtensionType->getIntParameter(1), ExtensionType->getIntParameter(2), in getImageType()
2732 ExtensionType->getIntParameter(3), ExtensionType->getIntParameter(4), in getImageType()
2733 SPIRV::ImageFormat::ImageFormat(ExtensionType->getIntParameter(5)), in getImageType()
2737 ExtensionType->getIntParameter(6))); in getImageType()
2746 return GR->getOrCreateOpTypeSampledImage(OpaqueImageType, MIRBuilder); in getSampledImageType()
2754 // Pointers-to-opaque-structs representing OpenCL types are first translated in parseBuiltinTypeNameToTargetExtType()
2755 // to equivalent SPIR-V types. OpenCL builtin type names should have the in parseBuiltinTypeNameToTargetExtType()
2763 NameWithParameters = OCLTypeRecord->SpirvTypeLiteral; in parseBuiltinTypeNameToTargetExtType()
2764 // Continue with the SPIR-V builtin type... in parseBuiltinTypeNameToTargetExtType()
2767 // Names of the opaque structs representing a SPIR-V builtins without in parseBuiltinTypeNameToTargetExtType()
2772 // Parameterized SPIR-V builtins names follow this format: in parseBuiltinTypeNameToTargetExtType()
2778 unsigned BaseNameLength = NameWithParameters.find('_') - 1; in parseBuiltinTypeNameToTargetExtType()
2791 "Invalid format of SPIR-V builtin parameter literal!"); in parseBuiltinTypeNameToTargetExtType()
2803 // In LLVM IR, SPIR-V and OpenCL builtin types are represented as either in lowerBuiltinType()
2804 // target(...) target extension types or pointers-to-opaque-structs. The in lowerBuiltinType()
2805 // approach relying on structs is deprecated and works only in the non-opaque in lowerBuiltinType()
2806 // pointer mode (-opaque-pointers=0). in lowerBuiltinType()
2808 // of Clang and LLVM/SPIR-V Translator, the pointers-to-opaque-structs are in lowerBuiltinType()
2814 OpaqueType->getStructName().str(), MIRBuilder.getContext()); in lowerBuiltinType()
2816 unsigned NumStartingVRegs = MIRBuilder.getMRI()->getNumVirtRegs(); in lowerBuiltinType()
2818 const StringRef Name = BuiltinType->getName(); in lowerBuiltinType()
2831 switch (TypeRecord->Opcode) { in lowerBuiltinType()
2839 TargetType = GR->getOrCreateOpTypeDeviceEvent(MIRBuilder); in lowerBuiltinType()
2858 if (NumStartingVRegs < MIRBuilder.getMRI()->getNumVirtRegs()) in lowerBuiltinType()
2859 buildOpName(GR->getSPIRVTypeID(TargetType), Name, MIRBuilder); in lowerBuiltinType()