shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity
llvm-project/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
Nathan Gauër 8cfda79105
[HLSL][SPIR-V] Implement vk::push_constant (#166793) 
Implements initial support for vk::push_constant.
As is, this allows handling simple push constants, but has one
main issue: layout can be incorrect (See #168401). The layout
issue being not only push-constant related, it's ignored for this PR.

The frontend part of the implementation is straightforward:
 - adding a new attribute
 - when targeting vulkan/spirv, we process it
 - global variables with this attribute gets a new AS:
   hlsl_push_constant

The IR has nothing specific, only some RO globals in this new AS.

On the SPIR-V side, we not convert this AS into a PushConstant storage
class. But this creates some issues: the variables in this storage class
must have a specific set of decoration to define their layout.

Current infra to create the SPIR-V types lacks the context required to
make this decision: no indication on the AS or context around the type
being created. Refactoring this would be a heavy task as it would
require getting this information in every place using the GR for type
creation.

Instead, we do something similar to CBuffers:
 - find all globals with this address space, and change their type to
   a target-specific type.
 - insert a new intrinsic in place of every reference to this global
   variable.

This allow the backend to handle both layout variables loads and type
lowering independently.

Type lowering has nothing specific: when we encounter a target extension
type with spirv.PushConstant, we lower this to the correct SPIR-V type
with the proper offset & block decorations.

As for the intrinsic, it's mostly a no-op, but required since we have
this target-specific type.

Note: this implementation prevents the static declaration of multiple
push constants in a single shader module. The actual specification is
more relaxed: there can be only one **used** push constant block per
entrypoint. To correctly implement this, we'd require to keep some
additional state to determine the list of statically used resources per
entrypoint. This shall be addressed as a follow-up (see #170310)
2025-12-18 11:01:11 +01:00

5067 lines
209 KiB
C++
Raw Blame History

//===- SPIRVInstructionSelector.cpp ------------------------------*- C++ -*-==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
//
// This file implements the targeting of the InstructionSelector class for
// SPIRV.
// TODO: This should be generated by TableGen.
//
//===----------------------------------------------------------------------===//
#include "MCTargetDesc/SPIRVBaseInfo.h"
#include "MCTargetDesc/SPIRVMCTargetDesc.h"
#include "SPIRV.h"
#include "SPIRVGlobalRegistry.h"
#include "SPIRVInstrInfo.h"
#include "SPIRVRegisterInfo.h"
#include "SPIRVTargetMachine.h"
#include "SPIRVUtils.h"
#include "llvm/ADT/APFloat.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/CodeGen/GlobalISel/GIMatchTableExecutorImpl.h"
#include "llvm/CodeGen/GlobalISel/GenericMachineInstrs.h"
#include "llvm/CodeGen/GlobalISel/InstructionSelector.h"
#include "llvm/CodeGen/MachineInstrBuilder.h"
#include "llvm/CodeGen/MachineRegisterInfo.h"
#include "llvm/CodeGen/Register.h"
#include "llvm/CodeGen/TargetOpcodes.h"
#include "llvm/IR/IntrinsicsSPIRV.h"
#include "llvm/Support/Debug.h"
#include "llvm/Support/ErrorHandling.h"
#define DEBUG_TYPE "spirv-isel"
using namespace llvm;
namespace CL = SPIRV::OpenCLExtInst;
namespace GL = SPIRV::GLSLExtInst;
using ExtInstList =
std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
namespace {
llvm::SPIRV::SelectionControl::SelectionControl
getSelectionOperandForImm(int Imm) {
if (Imm == 2)
return SPIRV::SelectionControl::Flatten;
if (Imm == 1)
return SPIRV::SelectionControl::DontFlatten;
if (Imm == 0)
return SPIRV::SelectionControl::None;
llvm_unreachable("Invalid immediate");
}
#define GET_GLOBALISEL_PREDICATE_BITSET
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_PREDICATE_BITSET
class SPIRVInstructionSelector : public InstructionSelector {
const SPIRVSubtarget &STI;
const SPIRVInstrInfo &TII;
const SPIRVRegisterInfo &TRI;
const RegisterBankInfo &RBI;
SPIRVGlobalRegistry &GR;
MachineRegisterInfo *MRI;
MachineFunction *HasVRegsReset = nullptr;
/// We need to keep track of the number we give to anonymous global values to
/// generate the same name every time when this is needed.
mutable DenseMap<const GlobalValue *, unsigned> UnnamedGlobalIDs;
SmallPtrSet<MachineInstr *, 8> DeadMIs;
public:
SPIRVInstructionSelector(const SPIRVTargetMachine &TM,
const SPIRVSubtarget &ST,
const RegisterBankInfo &RBI);
void setupMF(MachineFunction &MF, GISelValueTracking *VT,
CodeGenCoverage *CoverageInfo, ProfileSummaryInfo *PSI,
BlockFrequencyInfo *BFI) override;
// Common selection code. Instruction-specific selection occurs in spvSelect.
bool select(MachineInstr &I) override;
static const char *getName() { return DEBUG_TYPE; }
#define GET_GLOBALISEL_PREDICATES_DECL
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_PREDICATES_DECL
#define GET_GLOBALISEL_TEMPORARIES_DECL
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_TEMPORARIES_DECL
private:
void resetVRegsType(MachineFunction &MF);
void removeDeadInstruction(MachineInstr &MI) const;
void removeOpNamesForDeadMI(MachineInstr &MI) const;
// tblgen-erated 'select' implementation, used as the initial selector for
// the patterns that don't require complex C++.
bool selectImpl(MachineInstr &I, CodeGenCoverage &CoverageInfo) const;
// All instruction-specific selection that didn't happen in "select()".
// Is basically a large Switch/Case delegating to all other select method.
bool spvSelect(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectFirstBitHigh(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, bool IsSigned) const;
bool selectFirstBitLow(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectFirstBitSet16(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, unsigned ExtendOpcode,
unsigned BitSetOpcode) const;
bool selectFirstBitSet32(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, Register SrcReg,
unsigned BitSetOpcode) const;
bool selectFirstBitSet64(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, Register SrcReg,
unsigned BitSetOpcode, bool SwapPrimarySide) const;
bool selectFirstBitSet64Overflow(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, Register SrcReg,
unsigned BitSetOpcode,
bool SwapPrimarySide) const;
bool selectGlobalValue(Register ResVReg, MachineInstr &I,
const MachineInstr *Init = nullptr) const;
bool selectOpWithSrcs(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, std::vector<Register> SrcRegs,
unsigned Opcode) const;
bool selectUnOp(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
unsigned Opcode) const;
bool selectBitcast(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectLoad(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectStore(MachineInstr &I) const;
bool selectStackSave(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectStackRestore(MachineInstr &I) const;
bool selectMemOperation(Register ResVReg, MachineInstr &I) const;
Register getOrCreateMemSetGlobal(MachineInstr &I) const;
bool selectCopyMemory(MachineInstr &I, Register SrcReg) const;
bool selectCopyMemorySized(MachineInstr &I, Register SrcReg) const;
bool selectAtomicRMW(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, unsigned NewOpcode,
unsigned NegateOpcode = 0) const;
bool selectAtomicCmpXchg(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectFence(MachineInstr &I) const;
bool selectAddrSpaceCast(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectAnyOrAll(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, unsigned OpType) const;
bool selectAll(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectAny(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectBitreverse(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectBuildVector(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectSplatVector(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectCmp(Register ResVReg, const SPIRVType *ResType,
unsigned comparisonOpcode, MachineInstr &I) const;
bool selectDiscard(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectICmp(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectFCmp(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectSign(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectFloatDot(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectOverflowArith(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, unsigned Opcode) const;
bool selectDebugTrap(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectIntegerDot(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, bool Signed) const;
bool selectIntegerDotExpansion(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectOpIsInf(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectOpIsNan(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
template <bool Signed>
bool selectDot4AddPacked(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
template <bool Signed>
bool selectDot4AddPackedExpansion(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectWaveReduceMax(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, bool IsUnsigned) const;
bool selectWaveReduceMin(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, bool IsUnsigned) const;
bool selectWaveReduceSum(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectConst(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectSelect(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectSelectDefaultArgs(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, bool IsSigned) const;
bool selectIToF(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
bool IsSigned, unsigned Opcode) const;
bool selectExt(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
bool IsSigned) const;
bool selectTrunc(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectSUCmp(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
bool IsSigned) const;
bool selectIntToBool(Register IntReg, Register ResVReg, MachineInstr &I,
const SPIRVType *intTy, const SPIRVType *boolTy) const;
bool selectOpUndef(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectFreeze(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectIntrinsic(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectExtractVal(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectInsertVal(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectExtractElt(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectInsertElt(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectGEP(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectFrameIndex(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectAllocaArray(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectBranch(MachineInstr &I) const;
bool selectBranchCond(MachineInstr &I) const;
bool selectPhi(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectExtInst(Register ResVReg, const SPIRVType *RestType,
MachineInstr &I, GL::GLSLExtInst GLInst) const;
bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, CL::OpenCLExtInst CLInst) const;
bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, CL::OpenCLExtInst CLInst,
GL::GLSLExtInst GLInst) const;
bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, const ExtInstList &ExtInsts) const;
bool selectExtInstForLRound(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, CL::OpenCLExtInst CLInst,
GL::GLSLExtInst GLInst) const;
bool selectExtInstForLRound(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I,
const ExtInstList &ExtInsts) const;
bool selectLog10(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectSaturate(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectWaveOpInst(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, unsigned Opcode) const;
bool selectWaveActiveCountBits(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectUnmergeValues(MachineInstr &I) const;
bool selectHandleFromBinding(Register &ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectCounterHandleFromBinding(Register &ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const;
bool selectReadImageIntrinsic(Register &ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectImageWriteIntrinsic(MachineInstr &I) const;
bool selectResourceGetPointer(Register &ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectPushConstantGetPointer(Register &ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectResourceNonUniformIndex(Register &ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const;
bool selectModf(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectUpdateCounter(Register &ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectFrexp(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool selectDerivativeInst(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I, const unsigned DPdOpCode) const;
// Utilities
std::pair<Register, bool>
buildI32Constant(uint32_t Val, MachineInstr &I,
const SPIRVType *ResType = nullptr) const;
Register buildZerosVal(const SPIRVType *ResType, MachineInstr &I) const;
Register buildZerosValF(const SPIRVType *ResType, MachineInstr &I) const;
Register buildOnesVal(bool AllOnes, const SPIRVType *ResType,
MachineInstr &I) const;
Register buildOnesValF(const SPIRVType *ResType, MachineInstr &I) const;
bool wrapIntoSpecConstantOp(MachineInstr &I,
SmallVector<Register> &CompositeArgs) const;
Register getUcharPtrTypeReg(MachineInstr &I,
SPIRV::StorageClass::StorageClass SC) const;
MachineInstrBuilder buildSpecConstantOp(MachineInstr &I, Register Dest,
Register Src, Register DestType,
uint32_t Opcode) const;
MachineInstrBuilder buildConstGenericPtr(MachineInstr &I, Register SrcPtr,
SPIRVType *SrcPtrTy) const;
Register buildPointerToResource(const SPIRVType *ResType,
SPIRV::StorageClass::StorageClass SC,
uint32_t Set, uint32_t Binding,
uint32_t ArraySize, Register IndexReg,
StringRef Name,
MachineIRBuilder MIRBuilder) const;
SPIRVType *widenTypeToVec4(const SPIRVType *Type, MachineInstr &I) const;
bool extractSubvector(Register &ResVReg, const SPIRVType *ResType,
Register &ReadReg, MachineInstr &InsertionPoint) const;
bool generateImageReadOrFetch(Register &ResVReg, const SPIRVType *ResType,
Register ImageReg, Register IdxReg,
DebugLoc Loc, MachineInstr &Pos) const;
bool BuildCOPY(Register DestReg, Register SrcReg, MachineInstr &I) const;
bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
bool loadHandleBeforePosition(Register &HandleReg, const SPIRVType *ResType,
GIntrinsic &HandleDef, MachineInstr &Pos) const;
void decorateUsesAsNonUniform(Register &NonUniformReg) const;
void errorIfInstrOutsideShader(MachineInstr &I) const;
};
bool sampledTypeIsSignedInteger(const llvm::Type *HandleType) {
const TargetExtType *TET = cast<TargetExtType>(HandleType);
if (TET->getTargetExtName() == "spirv.Image") {
return false;
}
assert(TET->getTargetExtName() == "spirv.SignedImage");
return TET->getTypeParameter(0)->isIntegerTy();
}
} // end anonymous namespace
#define GET_GLOBALISEL_IMPL
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_IMPL
SPIRVInstructionSelector::SPIRVInstructionSelector(const SPIRVTargetMachine &TM,
const SPIRVSubtarget &ST,
const RegisterBankInfo &RBI)
: InstructionSelector(), STI(ST), TII(*ST.getInstrInfo()),
TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
MRI(nullptr),
#define GET_GLOBALISEL_PREDICATES_INIT
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_TEMPORARIES_INIT
{
}
void SPIRVInstructionSelector::setupMF(MachineFunction &MF,
GISelValueTracking *VT,
CodeGenCoverage *CoverageInfo,
ProfileSummaryInfo *PSI,
BlockFrequencyInfo *BFI) {
MRI = &MF.getRegInfo();
GR.setCurrentFunc(MF);
InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
}
// Ensure that register classes correspond to pattern matching rules.
void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
if (HasVRegsReset == &MF)
return;
HasVRegsReset = &MF;
MachineRegisterInfo &MRI = MF.getRegInfo();
for (unsigned I = 0, E = MRI.getNumVirtRegs(); I != E; ++I) {
Register Reg = Register::index2VirtReg(I);
LLT RegType = MRI.getType(Reg);
if (RegType.isScalar())
MRI.setType(Reg, LLT::scalar(64));
else if (RegType.isPointer())
MRI.setType(Reg, LLT::pointer(0, 64));
else if (RegType.isVector())
MRI.setType(Reg, LLT::fixed_vector(2, LLT::scalar(64)));
}
for (const auto &MBB : MF) {
for (const auto &MI : MBB) {
if (isPreISelGenericOpcode(MI.getOpcode()))
GR.erase(&MI);
if (MI.getOpcode() != SPIRV::ASSIGN_TYPE)
continue;
Register DstReg = MI.getOperand(0).getReg();
LLT DstType = MRI.getType(DstReg);
Register SrcReg = MI.getOperand(1).getReg();
LLT SrcType = MRI.getType(SrcReg);
if (DstType != SrcType)
MRI.setType(DstReg, MRI.getType(SrcReg));
const TargetRegisterClass *DstRC = MRI.getRegClassOrNull(DstReg);
const TargetRegisterClass *SrcRC = MRI.getRegClassOrNull(SrcReg);
if (DstRC != SrcRC && SrcRC)
MRI.setRegClass(DstReg, SrcRC);
}
}
}
// Return true if the type represents a constant register
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef,
SmallPtrSet<SPIRVType *, 4> &Visited) {
OpDef = passCopy(OpDef, MRI);
if (Visited.contains(OpDef))
return true;
Visited.insert(OpDef);
unsigned Opcode = OpDef->getOpcode();
switch (Opcode) {
case TargetOpcode::G_CONSTANT:
case TargetOpcode::G_FCONSTANT:
case TargetOpcode::G_IMPLICIT_DEF:
return true;
case TargetOpcode::G_INTRINSIC:
case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
return cast<GIntrinsic>(*OpDef).getIntrinsicID() ==
Intrinsic::spv_const_composite;
case TargetOpcode::G_BUILD_VECTOR:
case TargetOpcode::G_SPLAT_VECTOR: {
for (unsigned i = OpDef->getNumExplicitDefs(); i < OpDef->getNumOperands();
i++) {
MachineInstr *OpNestedDef =
OpDef->getOperand(i).isReg()
? MRI->getVRegDef(OpDef->getOperand(i).getReg())
: nullptr;
if (OpNestedDef && !isConstReg(MRI, OpNestedDef, Visited))
return false;
}
return true;
case SPIRV::OpConstantTrue:
case SPIRV::OpConstantFalse:
case SPIRV::OpConstantI:
case SPIRV::OpConstantF:
case SPIRV::OpConstantComposite:
case SPIRV::OpConstantCompositeContinuedINTEL:
case SPIRV::OpConstantSampler:
case SPIRV::OpConstantNull:
case SPIRV::OpUndef:
case SPIRV::OpConstantFunctionPointerINTEL:
return true;
}
}
return false;
}
// Return true if the virtual register represents a constant
static bool isConstReg(MachineRegisterInfo *MRI, Register OpReg) {
SmallPtrSet<SPIRVType *, 4> Visited;
if (MachineInstr *OpDef = MRI->getVRegDef(OpReg))
return isConstReg(MRI, OpDef, Visited);
return false;
}
// TODO(168736): We should make this either a flag in tabelgen
// or reduce our dependence on the global registry, so we can remove this
// function. It can easily be missed when new intrinsics are added.
// Most SPIR-V instrinsics are considered to have side-effects in their tablegen
// definition because they are referenced in the global registry. This is a list
// of intrinsics that have no side effects other than their references in the
// global registry.
static bool intrinsicHasSideEffects(Intrinsic::ID ID) {
switch (ID) {
// This is not an exhaustive list and may need to be updated.
case Intrinsic::spv_all:
case Intrinsic::spv_alloca:
case Intrinsic::spv_any:
case Intrinsic::spv_bitcast:
case Intrinsic::spv_const_composite:
case Intrinsic::spv_cross:
case Intrinsic::spv_degrees:
case Intrinsic::spv_distance:
case Intrinsic::spv_extractelt:
case Intrinsic::spv_extractv:
case Intrinsic::spv_faceforward:
case Intrinsic::spv_fdot:
case Intrinsic::spv_firstbitlow:
case Intrinsic::spv_firstbitshigh:
case Intrinsic::spv_firstbituhigh:
case Intrinsic::spv_frac:
case Intrinsic::spv_gep:
case Intrinsic::spv_global_offset:
case Intrinsic::spv_global_size:
case Intrinsic::spv_group_id:
case Intrinsic::spv_insertelt:
case Intrinsic::spv_insertv:
case Intrinsic::spv_isinf:
case Intrinsic::spv_isnan:
case Intrinsic::spv_lerp:
case Intrinsic::spv_length:
case Intrinsic::spv_normalize:
case Intrinsic::spv_num_subgroups:
case Intrinsic::spv_num_workgroups:
case Intrinsic::spv_ptrcast:
case Intrinsic::spv_radians:
case Intrinsic::spv_reflect:
case Intrinsic::spv_refract:
case Intrinsic::spv_resource_getpointer:
case Intrinsic::spv_resource_handlefrombinding:
case Intrinsic::spv_resource_handlefromimplicitbinding:
case Intrinsic::spv_resource_nonuniformindex:
case Intrinsic::spv_rsqrt:
case Intrinsic::spv_saturate:
case Intrinsic::spv_sdot:
case Intrinsic::spv_sign:
case Intrinsic::spv_smoothstep:
case Intrinsic::spv_step:
case Intrinsic::spv_subgroup_id:
case Intrinsic::spv_subgroup_local_invocation_id:
case Intrinsic::spv_subgroup_max_size:
case Intrinsic::spv_subgroup_size:
case Intrinsic::spv_thread_id:
case Intrinsic::spv_thread_id_in_group:
case Intrinsic::spv_udot:
case Intrinsic::spv_undef:
case Intrinsic::spv_value_md:
case Intrinsic::spv_workgroup_size:
return false;
default:
return true;
}
}
// TODO(168736): We should make this either a flag in tabelgen
// or reduce our dependence on the global registry, so we can remove this
// function. It can easily be missed when new intrinsics are added.
static bool isOpcodeWithNoSideEffects(unsigned Opcode) {
switch (Opcode) {
case SPIRV::OpTypeVoid:
case SPIRV::OpTypeBool:
case SPIRV::OpTypeInt:
case SPIRV::OpTypeFloat:
case SPIRV::OpTypeVector:
case SPIRV::OpTypeMatrix:
case SPIRV::OpTypeImage:
case SPIRV::OpTypeSampler:
case SPIRV::OpTypeSampledImage:
case SPIRV::OpTypeArray:
case SPIRV::OpTypeRuntimeArray:
case SPIRV::OpTypeStruct:
case SPIRV::OpTypeOpaque:
case SPIRV::OpTypePointer:
case SPIRV::OpTypeFunction:
case SPIRV::OpTypeEvent:
case SPIRV::OpTypeDeviceEvent:
case SPIRV::OpTypeReserveId:
case SPIRV::OpTypeQueue:
case SPIRV::OpTypePipe:
case SPIRV::OpTypeForwardPointer:
case SPIRV::OpTypePipeStorage:
case SPIRV::OpTypeNamedBarrier:
case SPIRV::OpTypeAccelerationStructureNV:
case SPIRV::OpTypeCooperativeMatrixNV:
case SPIRV::OpTypeCooperativeMatrixKHR:
return true;
default:
return false;
}
}
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI) {
// If there are no definitions, then assume there is some other
// side-effect that makes this instruction live.
if (MI.getNumDefs() == 0)
return false;
for (const auto &MO : MI.all_defs()) {
Register Reg = MO.getReg();
if (Reg.isPhysical()) {
LLVM_DEBUG(dbgs() << "Not dead: def of physical register " << Reg);
return false;
}
for (const auto &UseMI : MRI.use_nodbg_instructions(Reg)) {
if (UseMI.getOpcode() != SPIRV::OpName) {
LLVM_DEBUG(dbgs() << "Not dead: def " << MO << " has use in " << UseMI);
return false;
}
}
}
if (MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE || MI.isFakeUse() ||
MI.isLifetimeMarker()) {
LLVM_DEBUG(
dbgs()
<< "Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
return false;
}
if (MI.isPHI()) {
LLVM_DEBUG(dbgs() << "Dead: Phi instruction with no uses.\n");
return true;
}
// It is possible that the only side effect is that the instruction is
// referenced in the global registry. If that is the only side effect, the
// intrinsic is dead.
if (MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
const auto &Intr = cast<GIntrinsic>(MI);
if (!intrinsicHasSideEffects(Intr.getIntrinsicID())) {
LLVM_DEBUG(dbgs() << "Dead: Intrinsic with no real side effects.\n");
return true;
}
}
if (MI.mayStore() || MI.isCall() ||
(MI.mayLoad() && MI.hasOrderedMemoryRef()) || MI.isPosition() ||
MI.isDebugInstr() || MI.isTerminator() || MI.isJumpTableDebugInfo()) {
LLVM_DEBUG(dbgs() << "Not dead: instruction has side effects.\n");
return false;
}
if (isPreISelGenericOpcode(MI.getOpcode())) {
// TODO: Is there a generic way to check if the opcode has side effects?
LLVM_DEBUG(dbgs() << "Dead: Generic opcode with no uses.\n");
return true;
}
if (isOpcodeWithNoSideEffects(MI.getOpcode())) {
LLVM_DEBUG(dbgs() << "Dead: known opcode with no side effects\n");
return true;
}
return false;
}
void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &MI) const {
// Delete the OpName that uses the result if there is one.
for (const auto &MO : MI.all_defs()) {
Register Reg = MO.getReg();
if (Reg.isPhysical())
continue;
SmallVector<MachineInstr *, 4> UselessOpNames;
for (MachineInstr &UseMI : MRI->use_nodbg_instructions(Reg)) {
assert(UseMI.getOpcode() == SPIRV::OpName &&
"There is still a use of the dead function.");
UselessOpNames.push_back(&UseMI);
}
for (MachineInstr *OpNameMI : UselessOpNames) {
GR.invalidateMachineInstr(OpNameMI);
OpNameMI->eraseFromParent();
}
}
}
void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &MI) const {
salvageDebugInfo(*MRI, MI);
GR.invalidateMachineInstr(&MI);
removeOpNamesForDeadMI(MI);
MI.eraseFromParent();
}
bool SPIRVInstructionSelector::select(MachineInstr &I) {
resetVRegsType(*I.getParent()->getParent());
assert(I.getParent() && "Instruction should be in a basic block!");
assert(I.getParent()->getParent() && "Instruction should be in a function!");
LLVM_DEBUG(dbgs() << "Checking if instruction is dead: " << I;);
if (isDead(I, *MRI)) {
LLVM_DEBUG(dbgs() << "Instruction is dead.\n");
removeDeadInstruction(I);
return true;
}
Register Opcode = I.getOpcode();
// If it's not a GMIR instruction, we've selected it already.
if (!isPreISelGenericOpcode(Opcode)) {
if (Opcode == SPIRV::ASSIGN_TYPE) { // These pseudos aren't needed any more.
Register DstReg = I.getOperand(0).getReg();
Register SrcReg = I.getOperand(1).getReg();
auto *Def = MRI->getVRegDef(SrcReg);
if (isTypeFoldingSupported(Def->getOpcode()) &&
Def->getOpcode() != TargetOpcode::G_CONSTANT &&
Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
bool Res = false;
if (Def->getOpcode() == TargetOpcode::G_SELECT) {
Register SelectDstReg = Def->getOperand(0).getReg();
Res = selectSelect(SelectDstReg, GR.getSPIRVTypeForVReg(SelectDstReg),
*Def);
GR.invalidateMachineInstr(Def);
Def->removeFromParent();
MRI->replaceRegWith(DstReg, SelectDstReg);
GR.invalidateMachineInstr(&I);
I.removeFromParent();
} else
Res = selectImpl(I, *CoverageInfo);
LLVM_DEBUG({
if (!Res && Def->getOpcode() != TargetOpcode::G_CONSTANT) {
dbgs() << "Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
I.print(dbgs());
}
});
assert(Res || Def->getOpcode() == TargetOpcode::G_CONSTANT);
if (Res) {
if (!isTriviallyDead(*Def, *MRI) && isDead(*Def, *MRI))
DeadMIs.insert(Def);
return Res;
}
}
MRI->setRegClass(SrcReg, MRI->getRegClass(DstReg));
MRI->replaceRegWith(SrcReg, DstReg);
GR.invalidateMachineInstr(&I);
I.removeFromParent();
return true;
} else if (I.getNumDefs() == 1) {
// Make all vregs 64 bits (for SPIR-V IDs).
MRI->setType(I.getOperand(0).getReg(), LLT::scalar(64));
}
return constrainSelectedInstRegOperands(I, TII, TRI, RBI);
}
if (DeadMIs.contains(&I)) {
// if the instruction has been already made dead by folding it away
// erase it
LLVM_DEBUG(dbgs() << "Instruction is folded and dead.\n");
removeDeadInstruction(I);
return true;
}
if (I.getNumOperands() != I.getNumExplicitOperands()) {
LLVM_DEBUG(errs() << "Generic instr has unexpected implicit operands\n");
return false;
}
// Common code for getting return reg+type, and removing selected instr
// from parent occurs here. Instr-specific selection happens in spvSelect().
bool HasDefs = I.getNumDefs() > 0;
Register ResVReg = HasDefs ? I.getOperand(0).getReg() : Register(0);
SPIRVType *ResType = HasDefs ? GR.getSPIRVTypeForVReg(ResVReg) : nullptr;
assert(!HasDefs || ResType || I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
if (spvSelect(ResVReg, ResType, I)) {
if (HasDefs) // Make all vregs 64 bits (for SPIR-V IDs).
for (unsigned i = 0; i < I.getNumDefs(); ++i)
MRI->setType(I.getOperand(i).getReg(), LLT::scalar(64));
GR.invalidateMachineInstr(&I);
I.removeFromParent();
return true;
}
return false;
}
static bool mayApplyGenericSelection(unsigned Opcode) {
switch (Opcode) {
case TargetOpcode::G_CONSTANT:
case TargetOpcode::G_FCONSTANT:
return false;
case TargetOpcode::G_SADDO:
case TargetOpcode::G_SSUBO:
return true;
}
return isTypeFoldingSupported(Opcode);
}
bool SPIRVInstructionSelector::BuildCOPY(Register DestReg, Register SrcReg,
MachineInstr &I) const {
const TargetRegisterClass *DstRC = MRI->getRegClassOrNull(DestReg);
const TargetRegisterClass *SrcRC = MRI->getRegClassOrNull(SrcReg);
if (DstRC != SrcRC && SrcRC)
MRI->setRegClass(DestReg, SrcRC);
return BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
.addDef(DestReg)
.addUse(SrcReg)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::spvSelect(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
const unsigned Opcode = I.getOpcode();
if (mayApplyGenericSelection(Opcode))
return selectImpl(I, *CoverageInfo);
switch (Opcode) {
case TargetOpcode::G_CONSTANT:
case TargetOpcode::G_FCONSTANT:
return selectConst(ResVReg, ResType, I);
case TargetOpcode::G_GLOBAL_VALUE:
return selectGlobalValue(ResVReg, I);
case TargetOpcode::G_IMPLICIT_DEF:
return selectOpUndef(ResVReg, ResType, I);
case TargetOpcode::G_FREEZE:
return selectFreeze(ResVReg, ResType, I);
case TargetOpcode::G_INTRINSIC:
case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
case TargetOpcode::G_INTRINSIC_CONVERGENT:
case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
return selectIntrinsic(ResVReg, ResType, I);
case TargetOpcode::G_BITREVERSE:
return selectBitreverse(ResVReg, ResType, I);
case TargetOpcode::G_BUILD_VECTOR:
return selectBuildVector(ResVReg, ResType, I);
case TargetOpcode::G_SPLAT_VECTOR:
return selectSplatVector(ResVReg, ResType, I);
case TargetOpcode::G_SHUFFLE_VECTOR: {
MachineBasicBlock &BB = *I.getParent();
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorShuffle))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(1).getReg())
.addUse(I.getOperand(2).getReg());
for (auto V : I.getOperand(3).getShuffleMask())
MIB.addImm(V);
return MIB.constrainAllUses(TII, TRI, RBI);
}
case TargetOpcode::G_MEMMOVE:
case TargetOpcode::G_MEMCPY:
case TargetOpcode::G_MEMSET:
return selectMemOperation(ResVReg, I);
case TargetOpcode::G_ICMP:
return selectICmp(ResVReg, ResType, I);
case TargetOpcode::G_FCMP:
return selectFCmp(ResVReg, ResType, I);
case TargetOpcode::G_FRAME_INDEX:
return selectFrameIndex(ResVReg, ResType, I);
case TargetOpcode::G_LOAD:
return selectLoad(ResVReg, ResType, I);
case TargetOpcode::G_STORE:
return selectStore(I);
case TargetOpcode::G_BR:
return selectBranch(I);
case TargetOpcode::G_BRCOND:
return selectBranchCond(I);
case TargetOpcode::G_PHI:
return selectPhi(ResVReg, ResType, I);
case TargetOpcode::G_FPTOSI:
return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToS);
case TargetOpcode::G_FPTOUI:
return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToU);
case TargetOpcode::G_FPTOSI_SAT:
return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToS);
case TargetOpcode::G_FPTOUI_SAT:
return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToU);
case TargetOpcode::G_SITOFP:
return selectIToF(ResVReg, ResType, I, true, SPIRV::OpConvertSToF);
case TargetOpcode::G_UITOFP:
return selectIToF(ResVReg, ResType, I, false, SPIRV::OpConvertUToF);
case TargetOpcode::G_CTPOP:
return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitCount);
case TargetOpcode::G_SMIN:
return selectExtInst(ResVReg, ResType, I, CL::s_min, GL::SMin);
case TargetOpcode::G_UMIN:
return selectExtInst(ResVReg, ResType, I, CL::u_min, GL::UMin);
case TargetOpcode::G_SMAX:
return selectExtInst(ResVReg, ResType, I, CL::s_max, GL::SMax);
case TargetOpcode::G_UMAX:
return selectExtInst(ResVReg, ResType, I, CL::u_max, GL::UMax);
case TargetOpcode::G_SCMP:
return selectSUCmp(ResVReg, ResType, I, true);
case TargetOpcode::G_UCMP:
return selectSUCmp(ResVReg, ResType, I, false);
case TargetOpcode::G_LROUND:
case TargetOpcode::G_LLROUND: {
Register regForLround =
MRI->createVirtualRegister(MRI->getRegClass(ResVReg), "lround");
MRI->setRegClass(regForLround, &SPIRV::iIDRegClass);
GR.assignSPIRVTypeToVReg(GR.getSPIRVTypeForVReg(I.getOperand(1).getReg()),
regForLround, *(I.getParent()->getParent()));
selectExtInstForLRound(regForLround, GR.getSPIRVTypeForVReg(regForLround),
I, CL::round, GL::Round);
MachineBasicBlock &BB = *I.getParent();
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConvertFToS))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(regForLround);
return MIB.constrainAllUses(TII, TRI, RBI);
}
case TargetOpcode::G_STRICT_FMA:
case TargetOpcode::G_FMA:
return selectExtInst(ResVReg, ResType, I, CL::fma, GL::Fma);
case TargetOpcode::G_STRICT_FLDEXP:
return selectExtInst(ResVReg, ResType, I, CL::ldexp);
case TargetOpcode::G_FPOW:
return selectExtInst(ResVReg, ResType, I, CL::pow, GL::Pow);
case TargetOpcode::G_FPOWI:
return selectExtInst(ResVReg, ResType, I, CL::pown);
case TargetOpcode::G_FEXP:
return selectExtInst(ResVReg, ResType, I, CL::exp, GL::Exp);
case TargetOpcode::G_FEXP2:
return selectExtInst(ResVReg, ResType, I, CL::exp2, GL::Exp2);
case TargetOpcode::G_FMODF:
return selectModf(ResVReg, ResType, I);
case TargetOpcode::G_FLOG:
return selectExtInst(ResVReg, ResType, I, CL::log, GL::Log);
case TargetOpcode::G_FLOG2:
return selectExtInst(ResVReg, ResType, I, CL::log2, GL::Log2);
case TargetOpcode::G_FLOG10:
return selectLog10(ResVReg, ResType, I);
case TargetOpcode::G_FABS:
return selectExtInst(ResVReg, ResType, I, CL::fabs, GL::FAbs);
case TargetOpcode::G_ABS:
return selectExtInst(ResVReg, ResType, I, CL::s_abs, GL::SAbs);
case TargetOpcode::G_FMINNUM:
case TargetOpcode::G_FMINIMUM:
return selectExtInst(ResVReg, ResType, I, CL::fmin, GL::NMin);
case TargetOpcode::G_FMAXNUM:
case TargetOpcode::G_FMAXIMUM:
return selectExtInst(ResVReg, ResType, I, CL::fmax, GL::NMax);
case TargetOpcode::G_FCOPYSIGN:
return selectExtInst(ResVReg, ResType, I, CL::copysign);
case TargetOpcode::G_FCEIL:
return selectExtInst(ResVReg, ResType, I, CL::ceil, GL::Ceil);
case TargetOpcode::G_FFLOOR:
return selectExtInst(ResVReg, ResType, I, CL::floor, GL::Floor);
case TargetOpcode::G_FCOS:
return selectExtInst(ResVReg, ResType, I, CL::cos, GL::Cos);
case TargetOpcode::G_FSIN:
return selectExtInst(ResVReg, ResType, I, CL::sin, GL::Sin);
case TargetOpcode::G_FTAN:
return selectExtInst(ResVReg, ResType, I, CL::tan, GL::Tan);
case TargetOpcode::G_FACOS:
return selectExtInst(ResVReg, ResType, I, CL::acos, GL::Acos);
case TargetOpcode::G_FASIN:
return selectExtInst(ResVReg, ResType, I, CL::asin, GL::Asin);
case TargetOpcode::G_FATAN:
return selectExtInst(ResVReg, ResType, I, CL::atan, GL::Atan);
case TargetOpcode::G_FATAN2:
return selectExtInst(ResVReg, ResType, I, CL::atan2, GL::Atan2);
case TargetOpcode::G_FCOSH:
return selectExtInst(ResVReg, ResType, I, CL::cosh, GL::Cosh);
case TargetOpcode::G_FSINH:
return selectExtInst(ResVReg, ResType, I, CL::sinh, GL::Sinh);
case TargetOpcode::G_FTANH:
return selectExtInst(ResVReg, ResType, I, CL::tanh, GL::Tanh);
case TargetOpcode::G_STRICT_FSQRT:
case TargetOpcode::G_FSQRT:
return selectExtInst(ResVReg, ResType, I, CL::sqrt, GL::Sqrt);
case TargetOpcode::G_CTTZ:
case TargetOpcode::G_CTTZ_ZERO_UNDEF:
return selectExtInst(ResVReg, ResType, I, CL::ctz);
case TargetOpcode::G_CTLZ:
case TargetOpcode::G_CTLZ_ZERO_UNDEF:
return selectExtInst(ResVReg, ResType, I, CL::clz);
case TargetOpcode::G_INTRINSIC_ROUND:
return selectExtInst(ResVReg, ResType, I, CL::round, GL::Round);
case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
return selectExtInst(ResVReg, ResType, I, CL::rint, GL::RoundEven);
case TargetOpcode::G_INTRINSIC_TRUNC:
return selectExtInst(ResVReg, ResType, I, CL::trunc, GL::Trunc);
case TargetOpcode::G_FRINT:
case TargetOpcode::G_FNEARBYINT:
return selectExtInst(ResVReg, ResType, I, CL::rint, GL::RoundEven);
case TargetOpcode::G_SMULH:
return selectExtInst(ResVReg, ResType, I, CL::s_mul_hi);
case TargetOpcode::G_UMULH:
return selectExtInst(ResVReg, ResType, I, CL::u_mul_hi);
case TargetOpcode::G_SADDSAT:
return selectExtInst(ResVReg, ResType, I, CL::s_add_sat);
case TargetOpcode::G_UADDSAT:
return selectExtInst(ResVReg, ResType, I, CL::u_add_sat);
case TargetOpcode::G_SSUBSAT:
return selectExtInst(ResVReg, ResType, I, CL::s_sub_sat);
case TargetOpcode::G_USUBSAT:
return selectExtInst(ResVReg, ResType, I, CL::u_sub_sat);
case TargetOpcode::G_FFREXP:
return selectFrexp(ResVReg, ResType, I);
case TargetOpcode::G_UADDO:
return selectOverflowArith(ResVReg, ResType, I,
ResType->getOpcode() == SPIRV::OpTypeVector
? SPIRV::OpIAddCarryV
: SPIRV::OpIAddCarryS);
case TargetOpcode::G_USUBO:
return selectOverflowArith(ResVReg, ResType, I,
ResType->getOpcode() == SPIRV::OpTypeVector
? SPIRV::OpISubBorrowV
: SPIRV::OpISubBorrowS);
case TargetOpcode::G_UMULO:
return selectOverflowArith(ResVReg, ResType, I, SPIRV::OpUMulExtended);
case TargetOpcode::G_SMULO:
return selectOverflowArith(ResVReg, ResType, I, SPIRV::OpSMulExtended);
case TargetOpcode::G_SEXT:
return selectExt(ResVReg, ResType, I, true);
case TargetOpcode::G_ANYEXT:
case TargetOpcode::G_ZEXT:
return selectExt(ResVReg, ResType, I, false);
case TargetOpcode::G_TRUNC:
return selectTrunc(ResVReg, ResType, I);
case TargetOpcode::G_FPTRUNC:
case TargetOpcode::G_FPEXT:
return selectUnOp(ResVReg, ResType, I, SPIRV::OpFConvert);
case TargetOpcode::G_PTRTOINT:
return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertPtrToU);
case TargetOpcode::G_INTTOPTR:
return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertUToPtr);
case TargetOpcode::G_BITCAST:
return selectBitcast(ResVReg, ResType, I);
case TargetOpcode::G_ADDRSPACE_CAST:
return selectAddrSpaceCast(ResVReg, ResType, I);
case TargetOpcode::G_PTR_ADD: {
// Currently, we get G_PTR_ADD only applied to global variables.
assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());
Register GV = I.getOperand(1).getReg();
MachineRegisterInfo::def_instr_iterator II = MRI->def_instr_begin(GV);
(void)II;
assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
(*II).getOpcode() == TargetOpcode::COPY ||
(*II).getOpcode() == SPIRV::OpVariable) &&
getImm(I.getOperand(2), MRI));
// It may be the initialization of a global variable.
bool IsGVInit = false;
for (MachineRegisterInfo::use_instr_iterator
UseIt = MRI->use_instr_begin(I.getOperand(0).getReg()),
UseEnd = MRI->use_instr_end();
UseIt != UseEnd; UseIt = std::next(UseIt)) {
if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
(*UseIt).getOpcode() == SPIRV::OpVariable) {
IsGVInit = true;
break;
}
}
MachineBasicBlock &BB = *I.getParent();
if (!IsGVInit) {
SPIRVType *GVType = GR.getSPIRVTypeForVReg(GV);
SPIRVType *GVPointeeType = GR.getPointeeType(GVType);
SPIRVType *ResPointeeType = GR.getPointeeType(ResType);
if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
// Build a new virtual register that is associated with the required
// data type.
Register NewVReg = MRI->createGenericVirtualRegister(MRI->getType(GV));
MRI->setRegClass(NewVReg, MRI->getRegClass(GV));
// Having a correctly typed base we are ready to build the actually
// required GEP. It may not be a constant though, because all Operands
// of OpSpecConstantOp is to originate from other const instructions,
// and only the AccessChain named opcodes accept a global OpVariable
// instruction. We can't use an AccessChain opcode because of the type
// mismatch between result and base types.
if (!GR.isBitcastCompatible(ResType, GVType))
report_fatal_error(
"incompatible result and operand types in a bitcast");
Register ResTypeReg = GR.getSPIRVTypeID(ResType);
MachineInstrBuilder MIB =
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpBitcast))
.addDef(NewVReg)
.addUse(ResTypeReg)
.addUse(GV);
return MIB.constrainAllUses(TII, TRI, RBI) &&
BuildMI(BB, I, I.getDebugLoc(),
TII.get(STI.isLogicalSPIRV()
? SPIRV::OpInBoundsAccessChain
: SPIRV::OpInBoundsPtrAccessChain))
.addDef(ResVReg)
.addUse(ResTypeReg)
.addUse(NewVReg)
.addUse(I.getOperand(2).getReg())
.constrainAllUses(TII, TRI, RBI);
} else {
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addImm(
static_cast<uint32_t>(SPIRV::Opcode::InBoundsPtrAccessChain))
.addUse(GV)
.addUse(I.getOperand(2).getReg())
.constrainAllUses(TII, TRI, RBI);
}
}
// It's possible to translate G_PTR_ADD to OpSpecConstantOp: either to
// initialize a global variable with a constant expression (e.g., the test
// case opencl/basic/progvar_prog_scope_init.ll), or for another use case
Register Idx = buildZerosVal(GR.getOrCreateSPIRVIntegerType(32, I, TII), I);
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addImm(static_cast<uint32_t>(
SPIRV::Opcode::InBoundsPtrAccessChain))
.addUse(GV)
.addUse(Idx)
.addUse(I.getOperand(2).getReg());
return MIB.constrainAllUses(TII, TRI, RBI);
}
case TargetOpcode::G_ATOMICRMW_OR:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicOr);
case TargetOpcode::G_ATOMICRMW_ADD:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicIAdd);
case TargetOpcode::G_ATOMICRMW_AND:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicAnd);
case TargetOpcode::G_ATOMICRMW_MAX:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicSMax);
case TargetOpcode::G_ATOMICRMW_MIN:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicSMin);
case TargetOpcode::G_ATOMICRMW_SUB:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicISub);
case TargetOpcode::G_ATOMICRMW_XOR:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicXor);
case TargetOpcode::G_ATOMICRMW_UMAX:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicUMax);
case TargetOpcode::G_ATOMICRMW_UMIN:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicUMin);
case TargetOpcode::G_ATOMICRMW_XCHG:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicExchange);
case TargetOpcode::G_ATOMIC_CMPXCHG:
return selectAtomicCmpXchg(ResVReg, ResType, I);
case TargetOpcode::G_ATOMICRMW_FADD:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFAddEXT);
case TargetOpcode::G_ATOMICRMW_FSUB:
// Translate G_ATOMICRMW_FSUB to OpAtomicFAddEXT with negative value operand
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFAddEXT,
ResType->getOpcode() == SPIRV::OpTypeVector
? SPIRV::OpFNegateV
: SPIRV::OpFNegate);
case TargetOpcode::G_ATOMICRMW_FMIN:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFMinEXT);
case TargetOpcode::G_ATOMICRMW_FMAX:
return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFMaxEXT);
case TargetOpcode::G_FENCE:
return selectFence(I);
case TargetOpcode::G_STACKSAVE:
return selectStackSave(ResVReg, ResType, I);
case TargetOpcode::G_STACKRESTORE:
return selectStackRestore(I);
case TargetOpcode::G_UNMERGE_VALUES:
return selectUnmergeValues(I);
// Discard gen opcodes for intrinsics which we do not expect to actually
// represent code after lowering or intrinsics which are not implemented but
// should not crash when found in a customer's LLVM IR input.
case TargetOpcode::G_TRAP:
case TargetOpcode::G_UBSANTRAP:
case TargetOpcode::DBG_LABEL:
return true;
case TargetOpcode::G_DEBUGTRAP:
return selectDebugTrap(ResVReg, ResType, I);
default:
return false;
}
}
bool SPIRVInstructionSelector::selectDebugTrap(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
unsigned Opcode = SPIRV::OpNop;
MachineBasicBlock &BB = *I.getParent();
return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
GL::GLSLExtInst GLInst) const {
if (!STI.canUseExtInstSet(
SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
std::string DiagMsg;
raw_string_ostream OS(DiagMsg);
I.print(OS, true, false, false, false);
DiagMsg += " is only supported with the GLSL extended instruction set.\n";
report_fatal_error(DiagMsg.c_str(), false);
}
return selectExtInst(ResVReg, ResType, I,
{{SPIRV::InstructionSet::GLSL_std_450, GLInst}});
}
bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
CL::OpenCLExtInst CLInst) const {
return selectExtInst(ResVReg, ResType, I,
{{SPIRV::InstructionSet::OpenCL_std, CLInst}});
}
bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
CL::OpenCLExtInst CLInst,
GL::GLSLExtInst GLInst) const {
ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
{SPIRV::InstructionSet::GLSL_std_450, GLInst}};
return selectExtInst(ResVReg, ResType, I, ExtInsts);
}
bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
const ExtInstList &Insts) const {
for (const auto &Ex : Insts) {
SPIRV::InstructionSet::InstructionSet Set = Ex.first;
uint32_t Opcode = Ex.second;
if (STI.canUseExtInstSet(Set)) {
MachineBasicBlock &BB = *I.getParent();
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addImm(static_cast<uint32_t>(Set))
.addImm(Opcode)
.setMIFlags(I.getFlags());
const unsigned NumOps = I.getNumOperands();
unsigned Index = 1;
if (Index < NumOps &&
I.getOperand(Index).getType() ==
MachineOperand::MachineOperandType::MO_IntrinsicID)
Index = 2;
for (; Index < NumOps; ++Index)
MIB.add(I.getOperand(Index));
return MIB.constrainAllUses(TII, TRI, RBI);
}
}
return false;
}
bool SPIRVInstructionSelector::selectExtInstForLRound(
Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst) const {
ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
{SPIRV::InstructionSet::GLSL_std_450, GLInst}};
return selectExtInstForLRound(ResVReg, ResType, I, ExtInsts);
}
bool SPIRVInstructionSelector::selectExtInstForLRound(
Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
const ExtInstList &Insts) const {
for (const auto &Ex : Insts) {
SPIRV::InstructionSet::InstructionSet Set = Ex.first;
uint32_t Opcode = Ex.second;
if (STI.canUseExtInstSet(Set)) {
MachineBasicBlock &BB = *I.getParent();
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addImm(static_cast<uint32_t>(Set))
.addImm(Opcode);
const unsigned NumOps = I.getNumOperands();
unsigned Index = 1;
if (Index < NumOps &&
I.getOperand(Index).getType() ==
MachineOperand::MachineOperandType::MO_IntrinsicID)
Index = 2;
for (; Index < NumOps; ++Index)
MIB.add(I.getOperand(Index));
MIB.constrainAllUses(TII, TRI, RBI);
return true;
}
}
return false;
}
bool SPIRVInstructionSelector::selectFrexp(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
{SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
for (const auto &Ex : ExtInsts) {
SPIRV::InstructionSet::InstructionSet Set = Ex.first;
uint32_t Opcode = Ex.second;
if (!STI.canUseExtInstSet(Set))
continue;
MachineIRBuilder MIRBuilder(I);
SPIRVType *PointeeTy = GR.getSPIRVTypeForVReg(I.getOperand(1).getReg());
const SPIRVType *PointerType = GR.getOrCreateSPIRVPointerType(
PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
Register PointerVReg =
createVirtualRegister(PointerType, &GR, MRI, MRI->getMF());
auto It = getOpVariableMBBIt(I);
auto MIB = BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
.addDef(PointerVReg)
.addUse(GR.getSPIRVTypeID(PointerType))
.addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function))
.constrainAllUses(TII, TRI, RBI);
MIB = MIB &
BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addImm(static_cast<uint32_t>(Ex.first))
.addImm(Opcode)
.add(I.getOperand(2))
.addUse(PointerVReg)
.constrainAllUses(TII, TRI, RBI);
MIB = MIB &
BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
.addDef(I.getOperand(1).getReg())
.addUse(GR.getSPIRVTypeID(PointeeTy))
.addUse(PointerVReg)
.constrainAllUses(TII, TRI, RBI);
return MIB;
}
return false;
}
bool SPIRVInstructionSelector::selectOpWithSrcs(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
std::vector<Register> Srcs,
unsigned Opcode) const {
auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType));
for (Register SReg : Srcs) {
MIB.addUse(SReg);
}
return MIB.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectUnOp(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
unsigned Opcode) const {
if (STI.isPhysicalSPIRV() && I.getOperand(1).isReg()) {
Register SrcReg = I.getOperand(1).getReg();
bool IsGV = false;
for (MachineRegisterInfo::def_instr_iterator DefIt =
MRI->def_instr_begin(SrcReg);
DefIt != MRI->def_instr_end(); DefIt = std::next(DefIt)) {
unsigned DefOpCode = DefIt->getOpcode();
if (DefOpCode == SPIRV::ASSIGN_TYPE) {
// We need special handling to look through the type assignment and see
// if this is a constant or a global
if (auto *VRD = getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
DefOpCode = VRD->getOpcode();
}
if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
DefOpCode == TargetOpcode::G_CONSTANT ||
DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
IsGV = true;
break;
}
}
if (IsGV) {
uint32_t SpecOpcode = 0;
switch (Opcode) {
case SPIRV::OpConvertPtrToU:
SpecOpcode = static_cast<uint32_t>(SPIRV::Opcode::ConvertPtrToU);
break;
case SPIRV::OpConvertUToPtr:
SpecOpcode = static_cast<uint32_t>(SPIRV::Opcode::ConvertUToPtr);
break;
}
if (SpecOpcode)
return BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addImm(SpecOpcode)
.addUse(SrcReg)
.constrainAllUses(TII, TRI, RBI);
}
}
return selectOpWithSrcs(ResVReg, ResType, I, {I.getOperand(1).getReg()},
Opcode);
}
bool SPIRVInstructionSelector::selectBitcast(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
Register OpReg = I.getOperand(1).getReg();
SPIRVType *OpType = OpReg.isValid() ? GR.getSPIRVTypeForVReg(OpReg) : nullptr;
if (!GR.isBitcastCompatible(ResType, OpType))
report_fatal_error("incompatible result and operand types in a bitcast");
return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitcast);
}
static void addMemoryOperands(MachineMemOperand *MemOp,
MachineInstrBuilder &MIB,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry &GR) {
uint32_t SpvMemOp = static_cast<uint32_t>(SPIRV::MemoryOperand::None);
if (MemOp->isVolatile())
SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
if (MemOp->isNonTemporal())
SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
if (MemOp->getAlign().value())
SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
[[maybe_unused]] MachineInstr *AliasList = nullptr;
[[maybe_unused]] MachineInstr *NoAliasList = nullptr;
const SPIRVSubtarget *ST =
static_cast<const SPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());
if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
if (auto *MD = MemOp->getAAInfo().Scope) {
AliasList = GR.getOrAddMemAliasingINTELInst(MIRBuilder, MD);
if (AliasList)
SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
}
if (auto *MD = MemOp->getAAInfo().NoAlias) {
NoAliasList = GR.getOrAddMemAliasingINTELInst(MIRBuilder, MD);
if (NoAliasList)
SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
}
}
if (SpvMemOp != static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
MIB.addImm(SpvMemOp);
if (SpvMemOp & static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
MIB.addImm(MemOp->getAlign().value());
if (AliasList)
MIB.addUse(AliasList->getOperand(0).getReg());
if (NoAliasList)
MIB.addUse(NoAliasList->getOperand(0).getReg());
}
}
static void addMemoryOperands(uint64_t Flags, MachineInstrBuilder &MIB) {
uint32_t SpvMemOp = static_cast<uint32_t>(SPIRV::MemoryOperand::None);
if (Flags & MachineMemOperand::Flags::MOVolatile)
SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
if (Flags & MachineMemOperand::Flags::MONonTemporal)
SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
if (SpvMemOp != static_cast<uint32_t>(SPIRV::MemoryOperand::None))
MIB.addImm(SpvMemOp);
}
bool SPIRVInstructionSelector::selectLoad(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
unsigned OpOffset = isa<GIntrinsic>(I) ? 1 : 0;
Register Ptr = I.getOperand(1 + OpOffset).getReg();
auto *PtrDef = getVRegDef(*MRI, Ptr);
auto *IntPtrDef = dyn_cast<GIntrinsic>(PtrDef);
if (IntPtrDef &&
IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
Register HandleReg = IntPtrDef->getOperand(2).getReg();
SPIRVType *HandleType = GR.getSPIRVTypeForVReg(HandleReg);
if (HandleType->getOpcode() == SPIRV::OpTypeImage) {
Register NewHandleReg =
MRI->createVirtualRegister(MRI->getRegClass(HandleReg));
auto *HandleDef = cast<GIntrinsic>(getVRegDef(*MRI, HandleReg));
if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef, I)) {
return false;
}
Register IdxReg = IntPtrDef->getOperand(3).getReg();
return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
I.getDebugLoc(), I);
}
}
auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(Ptr);
if (!I.getNumMemOperands()) {
assert(I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
I.getOpcode() ==
TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
addMemoryOperands(I.getOperand(2 + OpOffset).getImm(), MIB);
} else {
MachineIRBuilder MIRBuilder(I);
addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
}
return MIB.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectStore(MachineInstr &I) const {
unsigned OpOffset = isa<GIntrinsic>(I) ? 1 : 0;
Register StoreVal = I.getOperand(0 + OpOffset).getReg();
Register Ptr = I.getOperand(1 + OpOffset).getReg();
auto *PtrDef = getVRegDef(*MRI, Ptr);
auto *IntPtrDef = dyn_cast<GIntrinsic>(PtrDef);
if (IntPtrDef &&
IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
Register HandleReg = IntPtrDef->getOperand(2).getReg();
Register NewHandleReg =
MRI->createVirtualRegister(MRI->getRegClass(HandleReg));
auto *HandleDef = cast<GIntrinsic>(getVRegDef(*MRI, HandleReg));
SPIRVType *HandleType = GR.getSPIRVTypeForVReg(HandleReg);
if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef, I)) {
return false;
}
Register IdxReg = IntPtrDef->getOperand(3).getReg();
if (HandleType->getOpcode() == SPIRV::OpTypeImage) {
auto BMI = BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
.addUse(NewHandleReg)
.addUse(IdxReg)
.addUse(StoreVal);
const llvm::Type *LLVMHandleType = GR.getTypeForSPIRVType(HandleType);
if (sampledTypeIsSignedInteger(LLVMHandleType))
BMI.addImm(0x1000); // SignExtend
return BMI.constrainAllUses(TII, TRI, RBI);
}
}
MachineBasicBlock &BB = *I.getParent();
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpStore))
.addUse(Ptr)
.addUse(StoreVal);
if (!I.getNumMemOperands()) {
assert(I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
I.getOpcode() ==
TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
addMemoryOperands(I.getOperand(2 + OpOffset).getImm(), MIB);
} else {
MachineIRBuilder MIRBuilder(I);
addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
}
return MIB.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectStackSave(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
report_fatal_error(
"llvm.stacksave intrinsic: this instruction requires the following "
"SPIR-V extension: SPV_INTEL_variable_length_array",
false);
MachineBasicBlock &BB = *I.getParent();
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSaveMemoryINTEL))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &I) const {
if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
report_fatal_error(
"llvm.stackrestore intrinsic: this instruction requires the following "
"SPIR-V extension: SPV_INTEL_variable_length_array",
false);
if (!I.getOperand(0).isReg())
return false;
MachineBasicBlock &BB = *I.getParent();
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpRestoreMemoryINTEL))
.addUse(I.getOperand(0).getReg())
.constrainAllUses(TII, TRI, RBI);
}
Register
SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &I) const {
MachineIRBuilder MIRBuilder(I);
assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());
// TODO: check if we have such GV, add init, use buildGlobalVariable.
unsigned Num = getIConstVal(I.getOperand(2).getReg(), MRI);
Function &CurFunction = GR.CurMF->getFunction();
Type *LLVMArrTy =
ArrayType::get(IntegerType::get(CurFunction.getContext(), 8), Num);
GlobalVariable *GV = new GlobalVariable(*CurFunction.getParent(), LLVMArrTy,
true, GlobalValue::InternalLinkage,
Constant::getNullValue(LLVMArrTy));
Type *ValTy = Type::getInt8Ty(I.getMF()->getFunction().getContext());
Type *ArrTy = ArrayType::get(ValTy, Num);
SPIRVType *VarTy = GR.getOrCreateSPIRVPointerType(
ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
SPIRVType *SpvArrTy = GR.getOrCreateSPIRVType(
ArrTy, MIRBuilder, SPIRV::AccessQualifier::None, false);
unsigned Val = getIConstVal(I.getOperand(1).getReg(), MRI);
Register Const = GR.getOrCreateConstIntArray(Val, Num, I, SpvArrTy, TII);
Register VarReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
auto MIBVar =
BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpVariable))
.addDef(VarReg)
.addUse(GR.getSPIRVTypeID(VarTy))
.addImm(SPIRV::StorageClass::UniformConstant)
.addUse(Const);
if (!MIBVar.constrainAllUses(TII, TRI, RBI))
return Register();
GR.add(GV, MIBVar);
GR.addGlobalObject(GV, GR.CurMF, VarReg);
buildOpDecorate(VarReg, I, TII, SPIRV::Decoration::Constant, {});
return VarReg;
}
bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &I,
Register SrcReg) const {
MachineBasicBlock &BB = *I.getParent();
Register DstReg = I.getOperand(0).getReg();
SPIRVType *DstTy = GR.getSPIRVTypeForVReg(DstReg);
SPIRVType *SrcTy = GR.getSPIRVTypeForVReg(SrcReg);
if (GR.getPointeeType(DstTy) != GR.getPointeeType(SrcTy))
report_fatal_error("OpCopyMemory requires operands to have the same type");
uint64_t CopySize = getIConstVal(I.getOperand(2).getReg(), MRI);
SPIRVType *PointeeTy = GR.getPointeeType(DstTy);
const Type *LLVMPointeeTy = GR.getTypeForSPIRVType(PointeeTy);
if (!LLVMPointeeTy)
report_fatal_error(
"Unable to determine pointee type size for OpCopyMemory");
const DataLayout &DL = I.getMF()->getFunction().getDataLayout();
if (CopySize != DL.getTypeStoreSize(const_cast<Type *>(LLVMPointeeTy)))
report_fatal_error(
"OpCopyMemory requires the size to match the pointee type size");
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCopyMemory))
.addUse(DstReg)
.addUse(SrcReg);
if (I.getNumMemOperands()) {
MachineIRBuilder MIRBuilder(I);
addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
}
return MIB.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &I,
Register SrcReg) const {
MachineBasicBlock &BB = *I.getParent();
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCopyMemorySized))
.addUse(I.getOperand(0).getReg())
.addUse(SrcReg)
.addUse(I.getOperand(2).getReg());
if (I.getNumMemOperands()) {
MachineIRBuilder MIRBuilder(I);
addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
}
return MIB.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectMemOperation(Register ResVReg,
MachineInstr &I) const {
Register SrcReg = I.getOperand(1).getReg();
bool Result = true;
if (I.getOpcode() == TargetOpcode::G_MEMSET) {
Register VarReg = getOrCreateMemSetGlobal(I);
if (!VarReg.isValid())
return false;
Type *ValTy = Type::getInt8Ty(I.getMF()->getFunction().getContext());
SPIRVType *SourceTy = GR.getOrCreateSPIRVPointerType(
ValTy, I, SPIRV::StorageClass::UniformConstant);
SrcReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
Result &= selectOpWithSrcs(SrcReg, SourceTy, I, {VarReg}, SPIRV::OpBitcast);
}
if (STI.isLogicalSPIRV()) {
Result &= selectCopyMemory(I, SrcReg);
} else {
Result &= selectCopyMemorySized(I, SrcReg);
}
if (ResVReg.isValid() && ResVReg != I.getOperand(0).getReg())
Result &= BuildCOPY(ResVReg, I.getOperand(0).getReg(), I);
return Result;
}
bool SPIRVInstructionSelector::selectAtomicRMW(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
unsigned NewOpcode,
unsigned NegateOpcode) const {
bool Result = true;
assert(I.hasOneMemOperand());
const MachineMemOperand *MemOp = *I.memoperands_begin();
uint32_t Scope = static_cast<uint32_t>(getMemScope(
GR.CurMF->getFunction().getContext(), MemOp->getSyncScopeID()));
auto ScopeConstant = buildI32Constant(Scope, I);
Register ScopeReg = ScopeConstant.first;
Result &= ScopeConstant.second;
Register Ptr = I.getOperand(1).getReg();
// TODO: Changed as it's implemented in the translator. See test/atomicrmw.ll
// auto ScSem =
// getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr));
AtomicOrdering AO = MemOp->getSuccessOrdering();
uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));
auto MemSemConstant = buildI32Constant(MemSem /*| ScSem*/, I);
Register MemSemReg = MemSemConstant.first;
Result &= MemSemConstant.second;
Register ValueReg = I.getOperand(2).getReg();
if (NegateOpcode != 0) {
// Translation with negative value operand is requested
Register TmpReg = createVirtualRegister(ResType, &GR, MRI, MRI->getMF());
Result &= selectOpWithSrcs(TmpReg, ResType, I, {ValueReg}, NegateOpcode);
ValueReg = TmpReg;
}
return Result &&
BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(NewOpcode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(Ptr)
.addUse(ScopeReg)
.addUse(MemSemReg)
.addUse(ValueReg)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &I) const {
unsigned ArgI = I.getNumOperands() - 1;
Register SrcReg =
I.getOperand(ArgI).isReg() ? I.getOperand(ArgI).getReg() : Register(0);
SPIRVType *SrcType =
SrcReg.isValid() ? GR.getSPIRVTypeForVReg(SrcReg) : nullptr;
if (!SrcType || SrcType->getOpcode() != SPIRV::OpTypeVector)
report_fatal_error(
"cannot select G_UNMERGE_VALUES with a non-vector argument");
SPIRVType *ScalarType =
GR.getSPIRVTypeForVReg(SrcType->getOperand(1).getReg());
MachineBasicBlock &BB = *I.getParent();
bool Res = false;
unsigned CurrentIndex = 0;
for (unsigned i = 0; i < I.getNumDefs(); ++i) {
Register ResVReg = I.getOperand(i).getReg();
SPIRVType *ResType = GR.getSPIRVTypeForVReg(ResVReg);
if (!ResType) {
LLT ResLLT = MRI->getType(ResVReg);
assert(ResLLT.isValid());
if (ResLLT.isVector()) {
ResType = GR.getOrCreateSPIRVVectorType(
ScalarType, ResLLT.getNumElements(), I, TII);
} else {
ResType = ScalarType;
}
MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
GR.assignSPIRVTypeToVReg(ResType, ResVReg, *GR.CurMF);
}
if (ResType->getOpcode() == SPIRV::OpTypeVector) {
Register UndefReg = GR.getOrCreateUndef(I, SrcType, TII);
auto MIB =
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorShuffle))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(SrcReg)
.addUse(UndefReg);
unsigned NumElements = GR.getScalarOrVectorComponentCount(ResType);
for (unsigned j = 0; j < NumElements; ++j) {
MIB.addImm(CurrentIndex + j);
}
CurrentIndex += NumElements;
Res |= MIB.constrainAllUses(TII, TRI, RBI);
} else {
auto MIB =
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(SrcReg)
.addImm(CurrentIndex);
CurrentIndex++;
Res |= MIB.constrainAllUses(TII, TRI, RBI);
}
}
return Res;
}
bool SPIRVInstructionSelector::selectFence(MachineInstr &I) const {
AtomicOrdering AO = AtomicOrdering(I.getOperand(0).getImm());
uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));
auto MemSemConstant = buildI32Constant(MemSem, I);
Register MemSemReg = MemSemConstant.first;
bool Result = MemSemConstant.second;
SyncScope::ID Ord = SyncScope::ID(I.getOperand(1).getImm());
uint32_t Scope = static_cast<uint32_t>(
getMemScope(GR.CurMF->getFunction().getContext(), Ord));
auto ScopeConstant = buildI32Constant(Scope, I);
Register ScopeReg = ScopeConstant.first;
Result &= ScopeConstant.second;
MachineBasicBlock &BB = *I.getParent();
return Result &&
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpMemoryBarrier))
.addUse(ScopeReg)
.addUse(MemSemReg)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectOverflowArith(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
unsigned Opcode) const {
Type *ResTy = nullptr;
StringRef ResName;
if (!GR.findValueAttrs(&I, ResTy, ResName))
report_fatal_error(
"Not enough info to select the arithmetic with overflow instruction");
if (!ResTy || !ResTy->isStructTy())
report_fatal_error("Expect struct type result for the arithmetic "
"with overflow instruction");
// "Result Type must be from OpTypeStruct. The struct must have two members,
// and the two members must be the same type."
Type *ResElemTy = cast<StructType>(ResTy)->getElementType(0);
ResTy = StructType::get(ResElemTy, ResElemTy);
// Build SPIR-V types and constant(s) if needed.
MachineIRBuilder MIRBuilder(I);
SPIRVType *StructType = GR.getOrCreateSPIRVType(
ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite, false);
assert(I.getNumDefs() > 1 && "Not enought operands");
SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
unsigned N = GR.getScalarOrVectorComponentCount(ResType);
if (N > 1)
BoolType = GR.getOrCreateSPIRVVectorType(BoolType, N, I, TII);
Register BoolTypeReg = GR.getSPIRVTypeID(BoolType);
Register ZeroReg = buildZerosVal(ResType, I);
// A new virtual register to store the result struct.
Register StructVReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
MRI->setRegClass(StructVReg, &SPIRV::IDRegClass);
// Build the result name if needed.
if (ResName.size() > 0)
buildOpName(StructVReg, ResName, MIRBuilder);
// Build the arithmetic with overflow instruction.
MachineBasicBlock &BB = *I.getParent();
auto MIB =
BuildMI(BB, MIRBuilder.getInsertPt(), I.getDebugLoc(), TII.get(Opcode))
.addDef(StructVReg)
.addUse(GR.getSPIRVTypeID(StructType));
for (unsigned i = I.getNumDefs(); i < I.getNumOperands(); ++i)
MIB.addUse(I.getOperand(i).getReg());
bool Result = MIB.constrainAllUses(TII, TRI, RBI);
// Build instructions to extract fields of the instruction's result.
// A new virtual register to store the higher part of the result struct.
Register HigherVReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
MRI->setRegClass(HigherVReg, &SPIRV::iIDRegClass);
for (unsigned i = 0; i < I.getNumDefs(); ++i) {
auto MIB =
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
.addDef(i == 1 ? HigherVReg : I.getOperand(i).getReg())
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(StructVReg)
.addImm(i);
Result &= MIB.constrainAllUses(TII, TRI, RBI);
}
// Build boolean value from the higher part.
return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpINotEqual))
.addDef(I.getOperand(1).getReg())
.addUse(BoolTypeReg)
.addUse(HigherVReg)
.addUse(ZeroReg)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectAtomicCmpXchg(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
bool Result = true;
Register ScopeReg;
Register MemSemEqReg;
Register MemSemNeqReg;
Register Ptr = I.getOperand(2).getReg();
if (!isa<GIntrinsic>(I)) {
assert(I.hasOneMemOperand());
const MachineMemOperand *MemOp = *I.memoperands_begin();
unsigned Scope = static_cast<uint32_t>(getMemScope(
GR.CurMF->getFunction().getContext(), MemOp->getSyncScopeID()));
auto ScopeConstant = buildI32Constant(Scope, I);
ScopeReg = ScopeConstant.first;
Result &= ScopeConstant.second;
unsigned ScSem = static_cast<uint32_t>(
getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr)));
AtomicOrdering AO = MemOp->getSuccessOrdering();
unsigned MemSemEq = static_cast<uint32_t>(getMemSemantics(AO)) | ScSem;
auto MemSemEqConstant = buildI32Constant(MemSemEq, I);
MemSemEqReg = MemSemEqConstant.first;
Result &= MemSemEqConstant.second;
AtomicOrdering FO = MemOp->getFailureOrdering();
unsigned MemSemNeq = static_cast<uint32_t>(getMemSemantics(FO)) | ScSem;
if (MemSemEq == MemSemNeq)
MemSemNeqReg = MemSemEqReg;
else {
auto MemSemNeqConstant = buildI32Constant(MemSemEq, I);
MemSemNeqReg = MemSemNeqConstant.first;
Result &= MemSemNeqConstant.second;
}
} else {
ScopeReg = I.getOperand(5).getReg();
MemSemEqReg = I.getOperand(6).getReg();
MemSemNeqReg = I.getOperand(7).getReg();
}
Register Cmp = I.getOperand(3).getReg();
Register Val = I.getOperand(4).getReg();
SPIRVType *SpvValTy = GR.getSPIRVTypeForVReg(Val);
Register ACmpRes = createVirtualRegister(SpvValTy, &GR, MRI, *I.getMF());
const DebugLoc &DL = I.getDebugLoc();
Result &=
BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpAtomicCompareExchange))
.addDef(ACmpRes)
.addUse(GR.getSPIRVTypeID(SpvValTy))
.addUse(Ptr)
.addUse(ScopeReg)
.addUse(MemSemEqReg)
.addUse(MemSemNeqReg)
.addUse(Val)
.addUse(Cmp)
.constrainAllUses(TII, TRI, RBI);
SPIRVType *BoolTy = GR.getOrCreateSPIRVBoolType(I, TII);
Register CmpSuccReg = createVirtualRegister(BoolTy, &GR, MRI, *I.getMF());
Result &= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpIEqual))
.addDef(CmpSuccReg)
.addUse(GR.getSPIRVTypeID(BoolTy))
.addUse(ACmpRes)
.addUse(Cmp)
.constrainAllUses(TII, TRI, RBI);
Register TmpReg = createVirtualRegister(ResType, &GR, MRI, *I.getMF());
Result &= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))
.addDef(TmpReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(ACmpRes)
.addUse(GR.getOrCreateUndef(I, ResType, TII))
.addImm(0)
.constrainAllUses(TII, TRI, RBI);
return Result &&
BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(CmpSuccReg)
.addUse(TmpReg)
.addImm(1)
.constrainAllUses(TII, TRI, RBI);
}
static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC) {
switch (SC) {
case SPIRV::StorageClass::DeviceOnlyINTEL:
case SPIRV::StorageClass::HostOnlyINTEL:
return true;
default:
return false;
}
}
// Returns true ResVReg is referred only from global vars and OpName's.
static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg) {
bool IsGRef = false;
bool IsAllowedRefs =
llvm::all_of(MRI->use_instructions(ResVReg), [&IsGRef](auto const &It) {
unsigned Opcode = It.getOpcode();
if (Opcode == SPIRV::OpConstantComposite ||
Opcode == SPIRV::OpVariable ||
isSpvIntrinsic(It, Intrinsic::spv_init_global))
return IsGRef = true;
return Opcode == SPIRV::OpName;
});
return IsAllowedRefs && IsGRef;
}
Register SPIRVInstructionSelector::getUcharPtrTypeReg(
MachineInstr &I, SPIRV::StorageClass::StorageClass SC) const {
return GR.getSPIRVTypeID(GR.getOrCreateSPIRVPointerType(
Type::getInt8Ty(I.getMF()->getFunction().getContext()), I, SC));
}
MachineInstrBuilder
SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &I, Register Dest,
Register Src, Register DestType,
uint32_t Opcode) const {
return BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
.addDef(Dest)
.addUse(DestType)
.addImm(Opcode)
.addUse(Src);
}
MachineInstrBuilder
SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &I, Register SrcPtr,
SPIRVType *SrcPtrTy) const {
SPIRVType *GenericPtrTy =
GR.changePointerStorageClass(SrcPtrTy, SPIRV::StorageClass::Generic, I);
Register Tmp = MRI->createVirtualRegister(&SPIRV::pIDRegClass);
MRI->setType(Tmp, LLT::pointer(storageClassToAddressSpace(
SPIRV::StorageClass::Generic),
GR.getPointerSize()));
MachineFunction *MF = I.getParent()->getParent();
GR.assignSPIRVTypeToVReg(GenericPtrTy, Tmp, *MF);
MachineInstrBuilder MIB = buildSpecConstantOp(
I, Tmp, SrcPtr, GR.getSPIRVTypeID(GenericPtrTy),
static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric));
GR.add(MIB.getInstr(), MIB);
return MIB;
}
// In SPIR-V address space casting can only happen to and from the Generic
// storage class. We can also only cast Workgroup, CrossWorkgroup, or Function
// pointers to and from Generic pointers. As such, we can convert e.g. from
// Workgroup to Function by going via a Generic pointer as an intermediary. All
// other combinations can only be done by a bitcast, and are probably not safe.
bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
MachineBasicBlock &BB = *I.getParent();
const DebugLoc &DL = I.getDebugLoc();
Register SrcPtr = I.getOperand(1).getReg();
SPIRVType *SrcPtrTy = GR.getSPIRVTypeForVReg(SrcPtr);
// don't generate a cast for a null that may be represented by OpTypeInt
if (SrcPtrTy->getOpcode() != SPIRV::OpTypePointer ||
ResType->getOpcode() != SPIRV::OpTypePointer)
return BuildCOPY(ResVReg, SrcPtr, I);
SPIRV::StorageClass::StorageClass SrcSC = GR.getPointerStorageClass(SrcPtrTy);
SPIRV::StorageClass::StorageClass DstSC = GR.getPointerStorageClass(ResType);
if (isASCastInGVar(MRI, ResVReg)) {
// AddrSpaceCast uses within OpVariable and OpConstantComposite instructions
// are expressed by OpSpecConstantOp with an Opcode.
// TODO: maybe insert a check whether the Kernel capability was declared and
// so PtrCastToGeneric/GenericCastToPtr are available.
unsigned SpecOpcode =
DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC)
? static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric)
: (SrcSC == SPIRV::StorageClass::Generic &&
isGenericCastablePtr(DstSC)
? static_cast<uint32_t>(SPIRV::Opcode::GenericCastToPtr)
: 0);
// TODO: OpConstantComposite expects i8*, so we are forced to forget a
// correct value of ResType and use general i8* instead. Maybe this should
// be addressed in the emit-intrinsic step to infer a correct
// OpConstantComposite type.
if (SpecOpcode) {
return buildSpecConstantOp(I, ResVReg, SrcPtr,
getUcharPtrTypeReg(I, DstSC), SpecOpcode)
.constrainAllUses(TII, TRI, RBI);
} else if (isGenericCastablePtr(SrcSC) && isGenericCastablePtr(DstSC)) {
MachineInstrBuilder MIB = buildConstGenericPtr(I, SrcPtr, SrcPtrTy);
return MIB.constrainAllUses(TII, TRI, RBI) &&
buildSpecConstantOp(
I, ResVReg, MIB->getOperand(0).getReg(),
getUcharPtrTypeReg(I, DstSC),
static_cast<uint32_t>(SPIRV::Opcode::GenericCastToPtr))
.constrainAllUses(TII, TRI, RBI);
}
}
// don't generate a cast between identical storage classes
if (SrcSC == DstSC)
return BuildCOPY(ResVReg, SrcPtr, I);
if ((SrcSC == SPIRV::StorageClass::Function &&
DstSC == SPIRV::StorageClass::Private) ||
(DstSC == SPIRV::StorageClass::Function &&
SrcSC == SPIRV::StorageClass::Private))
return BuildCOPY(ResVReg, SrcPtr, I);
// Casting from an eligible pointer to Generic.
if (DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC))
return selectUnOp(ResVReg, ResType, I, SPIRV::OpPtrCastToGeneric);
// Casting from Generic to an eligible pointer.
if (SrcSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(DstSC))
return selectUnOp(ResVReg, ResType, I, SPIRV::OpGenericCastToPtr);
// Casting between 2 eligible pointers using Generic as an intermediary.
if (isGenericCastablePtr(SrcSC) && isGenericCastablePtr(DstSC)) {
SPIRVType *GenericPtrTy =
GR.changePointerStorageClass(SrcPtrTy, SPIRV::StorageClass::Generic, I);
Register Tmp = createVirtualRegister(GenericPtrTy, &GR, MRI, MRI->getMF());
bool Result = BuildMI(BB, I, DL, TII.get(SPIRV::OpPtrCastToGeneric))
.addDef(Tmp)
.addUse(GR.getSPIRVTypeID(GenericPtrTy))
.addUse(SrcPtr)
.constrainAllUses(TII, TRI, RBI);
return Result && BuildMI(BB, I, DL, TII.get(SPIRV::OpGenericCastToPtr))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(Tmp)
.constrainAllUses(TII, TRI, RBI);
}
// Check if instructions from the SPV_INTEL_usm_storage_classes extension may
// be applied
if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::CrossWorkgroup)
return selectUnOp(ResVReg, ResType, I,
SPIRV::OpPtrCastToCrossWorkgroupINTEL);
if (SrcSC == SPIRV::StorageClass::CrossWorkgroup && isUSMStorageClass(DstSC))
return selectUnOp(ResVReg, ResType, I,
SPIRV::OpCrossWorkgroupCastToPtrINTEL);
if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::Generic)
return selectUnOp(ResVReg, ResType, I, SPIRV::OpPtrCastToGeneric);
if (SrcSC == SPIRV::StorageClass::Generic && isUSMStorageClass(DstSC))
return selectUnOp(ResVReg, ResType, I, SPIRV::OpGenericCastToPtr);
// Bitcast for pointers requires that the address spaces must match
return false;
}
static unsigned getFCmpOpcode(unsigned PredNum) {
auto Pred = static_cast<CmpInst::Predicate>(PredNum);
switch (Pred) {
case CmpInst::FCMP_OEQ:
return SPIRV::OpFOrdEqual;
case CmpInst::FCMP_OGE:
return SPIRV::OpFOrdGreaterThanEqual;
case CmpInst::FCMP_OGT:
return SPIRV::OpFOrdGreaterThan;
case CmpInst::FCMP_OLE:
return SPIRV::OpFOrdLessThanEqual;
case CmpInst::FCMP_OLT:
return SPIRV::OpFOrdLessThan;
case CmpInst::FCMP_ONE:
return SPIRV::OpFOrdNotEqual;
case CmpInst::FCMP_ORD:
return SPIRV::OpOrdered;
case CmpInst::FCMP_UEQ:
return SPIRV::OpFUnordEqual;
case CmpInst::FCMP_UGE:
return SPIRV::OpFUnordGreaterThanEqual;
case CmpInst::FCMP_UGT:
return SPIRV::OpFUnordGreaterThan;
case CmpInst::FCMP_ULE:
return SPIRV::OpFUnordLessThanEqual;
case CmpInst::FCMP_ULT:
return SPIRV::OpFUnordLessThan;
case CmpInst::FCMP_UNE:
return SPIRV::OpFUnordNotEqual;
case CmpInst::FCMP_UNO:
return SPIRV::OpUnordered;
default:
llvm_unreachable("Unknown predicate type for FCmp");
}
}
static unsigned getICmpOpcode(unsigned PredNum) {
auto Pred = static_cast<CmpInst::Predicate>(PredNum);
switch (Pred) {
case CmpInst::ICMP_EQ:
return SPIRV::OpIEqual;
case CmpInst::ICMP_NE:
return SPIRV::OpINotEqual;
case CmpInst::ICMP_SGE:
return SPIRV::OpSGreaterThanEqual;
case CmpInst::ICMP_SGT:
return SPIRV::OpSGreaterThan;
case CmpInst::ICMP_SLE:
return SPIRV::OpSLessThanEqual;
case CmpInst::ICMP_SLT:
return SPIRV::OpSLessThan;
case CmpInst::ICMP_UGE:
return SPIRV::OpUGreaterThanEqual;
case CmpInst::ICMP_UGT:
return SPIRV::OpUGreaterThan;
case CmpInst::ICMP_ULE:
return SPIRV::OpULessThanEqual;
case CmpInst::ICMP_ULT:
return SPIRV::OpULessThan;
default:
llvm_unreachable("Unknown predicate type for ICmp");
}
}
static unsigned getPtrCmpOpcode(unsigned Pred) {
switch (static_cast<CmpInst::Predicate>(Pred)) {
case CmpInst::ICMP_EQ:
return SPIRV::OpPtrEqual;
case CmpInst::ICMP_NE:
return SPIRV::OpPtrNotEqual;
default:
llvm_unreachable("Unknown predicate type for pointer comparison");
}
}
// Return the logical operation, or abort if none exists.
static unsigned getBoolCmpOpcode(unsigned PredNum) {
auto Pred = static_cast<CmpInst::Predicate>(PredNum);
switch (Pred) {
case CmpInst::ICMP_EQ:
return SPIRV::OpLogicalEqual;
case CmpInst::ICMP_NE:
return SPIRV::OpLogicalNotEqual;
default:
llvm_unreachable("Unknown predicate type for Bool comparison");
}
}
static APFloat getZeroFP(const Type *LLVMFloatTy) {
if (!LLVMFloatTy)
return APFloat::getZero(APFloat::IEEEsingle());
switch (LLVMFloatTy->getScalarType()->getTypeID()) {
case Type::HalfTyID:
return APFloat::getZero(APFloat::IEEEhalf());
default:
case Type::FloatTyID:
return APFloat::getZero(APFloat::IEEEsingle());
case Type::DoubleTyID:
return APFloat::getZero(APFloat::IEEEdouble());
}
}
static APFloat getOneFP(const Type *LLVMFloatTy) {
if (!LLVMFloatTy)
return APFloat::getOne(APFloat::IEEEsingle());
switch (LLVMFloatTy->getScalarType()->getTypeID()) {
case Type::HalfTyID:
return APFloat::getOne(APFloat::IEEEhalf());
default:
case Type::FloatTyID:
return APFloat::getOne(APFloat::IEEEsingle());
case Type::DoubleTyID:
return APFloat::getOne(APFloat::IEEEdouble());
}
}
bool SPIRVInstructionSelector::selectAnyOrAll(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
unsigned OpAnyOrAll) const {
assert(I.getNumOperands() == 3);
assert(I.getOperand(2).isReg());
MachineBasicBlock &BB = *I.getParent();
Register InputRegister = I.getOperand(2).getReg();
SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
if (!InputType)
report_fatal_error("Input Type could not be determined.");
bool IsBoolTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeBool);
bool IsVectorTy = InputType->getOpcode() == SPIRV::OpTypeVector;
if (IsBoolTy && !IsVectorTy) {
assert(ResVReg == I.getOperand(0).getReg());
return BuildCOPY(ResVReg, InputRegister, I);
}
bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
unsigned SpirvNotEqualId =
IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
SPIRVType *SpvBoolScalarTy = GR.getOrCreateSPIRVBoolType(I, TII);
SPIRVType *SpvBoolTy = SpvBoolScalarTy;
Register NotEqualReg = ResVReg;
if (IsVectorTy) {
NotEqualReg =
IsBoolTy ? InputRegister
: createVirtualRegister(SpvBoolTy, &GR, MRI, MRI->getMF());
const unsigned NumElts = InputType->getOperand(2).getImm();
SpvBoolTy = GR.getOrCreateSPIRVVectorType(SpvBoolTy, NumElts, I, TII);
}
bool Result = true;
if (!IsBoolTy) {
Register ConstZeroReg =
IsFloatTy ? buildZerosValF(InputType, I) : buildZerosVal(InputType, I);
Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SpirvNotEqualId))
.addDef(NotEqualReg)
.addUse(GR.getSPIRVTypeID(SpvBoolTy))
.addUse(InputRegister)
.addUse(ConstZeroReg)
.constrainAllUses(TII, TRI, RBI);
}
if (!IsVectorTy)
return Result;
return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(OpAnyOrAll))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(SpvBoolScalarTy))
.addUse(NotEqualReg)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectAll(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
return selectAnyOrAll(ResVReg, ResType, I, SPIRV::OpAll);
}
bool SPIRVInstructionSelector::selectAny(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
return selectAnyOrAll(ResVReg, ResType, I, SPIRV::OpAny);
}
// Select the OpDot instruction for the given float dot
bool SPIRVInstructionSelector::selectFloatDot(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
assert(I.getNumOperands() == 4);
assert(I.getOperand(2).isReg());
assert(I.getOperand(3).isReg());
[[maybe_unused]] SPIRVType *VecType =
GR.getSPIRVTypeForVReg(I.getOperand(2).getReg());
assert(VecType->getOpcode() == SPIRV::OpTypeVector &&
GR.getScalarOrVectorComponentCount(VecType) > 1 &&
"dot product requires a vector of at least 2 components");
[[maybe_unused]] SPIRVType *EltType =
GR.getSPIRVTypeForVReg(VecType->getOperand(1).getReg());
assert(EltType->getOpcode() == SPIRV::OpTypeFloat);
MachineBasicBlock &BB = *I.getParent();
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpDot))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(2).getReg())
.addUse(I.getOperand(3).getReg())
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectIntegerDot(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
bool Signed) const {
assert(I.getNumOperands() == 4);
assert(I.getOperand(2).isReg());
assert(I.getOperand(3).isReg());
MachineBasicBlock &BB = *I.getParent();
auto DotOp = Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
return BuildMI(BB, I, I.getDebugLoc(), TII.get(DotOp))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(2).getReg())
.addUse(I.getOperand(3).getReg())
.constrainAllUses(TII, TRI, RBI);
}
// Since pre-1.6 SPIRV has no integer dot implementation,
// expand by piecewise multiplying and adding the results
bool SPIRVInstructionSelector::selectIntegerDotExpansion(
Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
assert(I.getNumOperands() == 4);
assert(I.getOperand(2).isReg());
assert(I.getOperand(3).isReg());
MachineBasicBlock &BB = *I.getParent();
// Multiply the vectors, then sum the results
Register Vec0 = I.getOperand(2).getReg();
Register Vec1 = I.getOperand(3).getReg();
Register TmpVec = MRI->createVirtualRegister(GR.getRegClass(ResType));
SPIRVType *VecType = GR.getSPIRVTypeForVReg(Vec0);
bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIMulV))
.addDef(TmpVec)
.addUse(GR.getSPIRVTypeID(VecType))
.addUse(Vec0)
.addUse(Vec1)
.constrainAllUses(TII, TRI, RBI);
assert(VecType->getOpcode() == SPIRV::OpTypeVector &&
GR.getScalarOrVectorComponentCount(VecType) > 1 &&
"dot product requires a vector of at least 2 components");
Register Res = MRI->createVirtualRegister(GR.getRegClass(ResType));
Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
.addDef(Res)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(TmpVec)
.addImm(0)
.constrainAllUses(TII, TRI, RBI);
for (unsigned i = 1; i < GR.getScalarOrVectorComponentCount(VecType); i++) {
Register Elt = MRI->createVirtualRegister(GR.getRegClass(ResType));
Result &=
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
.addDef(Elt)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(TmpVec)
.addImm(i)
.constrainAllUses(TII, TRI, RBI);
Register Sum = i < GR.getScalarOrVectorComponentCount(VecType) - 1
? MRI->createVirtualRegister(GR.getRegClass(ResType))
: ResVReg;
Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
.addDef(Sum)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(Res)
.addUse(Elt)
.constrainAllUses(TII, TRI, RBI);
Res = Sum;
}
return Result;
}
bool SPIRVInstructionSelector::selectOpIsInf(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
MachineBasicBlock &BB = *I.getParent();
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIsInf))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(2).getReg())
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectOpIsNan(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
MachineBasicBlock &BB = *I.getParent();
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIsNan))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(2).getReg())
.constrainAllUses(TII, TRI, RBI);
}
template <bool Signed>
bool SPIRVInstructionSelector::selectDot4AddPacked(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
assert(I.getNumOperands() == 5);
assert(I.getOperand(2).isReg());
assert(I.getOperand(3).isReg());
assert(I.getOperand(4).isReg());
MachineBasicBlock &BB = *I.getParent();
Register Acc = I.getOperand(2).getReg();
Register X = I.getOperand(3).getReg();
Register Y = I.getOperand(4).getReg();
auto DotOp = Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
Register Dot = MRI->createVirtualRegister(GR.getRegClass(ResType));
bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(DotOp))
.addDef(Dot)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(X)
.addUse(Y)
.constrainAllUses(TII, TRI, RBI);
return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(Dot)
.addUse(Acc)
.constrainAllUses(TII, TRI, RBI);
}
// Since pre-1.6 SPIRV has no DotProductInput4x8BitPacked implementation,
// extract the elements of the packed inputs, multiply them and add the result
// to the accumulator.
template <bool Signed>
bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
assert(I.getNumOperands() == 5);
assert(I.getOperand(2).isReg());
assert(I.getOperand(3).isReg());
assert(I.getOperand(4).isReg());
MachineBasicBlock &BB = *I.getParent();
bool Result = true;
Register Acc = I.getOperand(2).getReg();
Register X = I.getOperand(3).getReg();
Register Y = I.getOperand(4).getReg();
SPIRVType *EltType = GR.getOrCreateSPIRVIntegerType(8, I, TII);
auto ExtractOp =
Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
bool ZeroAsNull = !STI.isShader();
// Extract the i8 element, multiply and add it to the accumulator
for (unsigned i = 0; i < 4; i++) {
// A[i]
Register AElt = MRI->createVirtualRegister(&SPIRV::IDRegClass);
Result &=
BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
.addDef(AElt)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(X)
.addUse(GR.getOrCreateConstInt(i * 8, I, EltType, TII, ZeroAsNull))
.addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
.constrainAllUses(TII, TRI, RBI);
// B[i]
Register BElt = MRI->createVirtualRegister(&SPIRV::IDRegClass);
Result &=
BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
.addDef(BElt)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(Y)
.addUse(GR.getOrCreateConstInt(i * 8, I, EltType, TII, ZeroAsNull))
.addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
.constrainAllUses(TII, TRI, RBI);
// A[i] * B[i]
Register Mul = MRI->createVirtualRegister(&SPIRV::IDRegClass);
Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIMulS))
.addDef(Mul)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(AElt)
.addUse(BElt)
.constrainAllUses(TII, TRI, RBI);
// Discard 24 highest-bits so that stored i32 register is i8 equivalent
Register MaskMul = MRI->createVirtualRegister(&SPIRV::IDRegClass);
Result &=
BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
.addDef(MaskMul)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(Mul)
.addUse(GR.getOrCreateConstInt(0, I, EltType, TII, ZeroAsNull))
.addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
.constrainAllUses(TII, TRI, RBI);
// Acc = Acc + A[i] * B[i]
Register Sum =
i < 3 ? MRI->createVirtualRegister(&SPIRV::IDRegClass) : ResVReg;
Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
.addDef(Sum)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(Acc)
.addUse(MaskMul)
.constrainAllUses(TII, TRI, RBI);
Acc = Sum;
}
return Result;
}
/// Transform saturate(x) to clamp(x, 0.0f, 1.0f) as SPIRV
/// does not have a saturate builtin.
bool SPIRVInstructionSelector::selectSaturate(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
assert(I.getNumOperands() == 3);
assert(I.getOperand(2).isReg());
MachineBasicBlock &BB = *I.getParent();
Register VZero = buildZerosValF(ResType, I);
Register VOne = buildOnesValF(ResType, I);
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
.addImm(GL::FClamp)
.addUse(I.getOperand(2).getReg())
.addUse(VZero)
.addUse(VOne)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectSign(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
assert(I.getNumOperands() == 3);
assert(I.getOperand(2).isReg());
MachineBasicBlock &BB = *I.getParent();
Register InputRegister = I.getOperand(2).getReg();
SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
auto &DL = I.getDebugLoc();
if (!InputType)
report_fatal_error("Input Type could not be determined.");
bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
unsigned SignBitWidth = GR.getScalarOrVectorBitWidth(InputType);
unsigned ResBitWidth = GR.getScalarOrVectorBitWidth(ResType);
bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
Register SignReg = NeedsConversion
? MRI->createVirtualRegister(&SPIRV::IDRegClass)
: ResVReg;
bool Result =
BuildMI(BB, I, DL, TII.get(SPIRV::OpExtInst))
.addDef(SignReg)
.addUse(GR.getSPIRVTypeID(InputType))
.addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
.addImm(SignOpcode)
.addUse(InputRegister)
.constrainAllUses(TII, TRI, RBI);
if (NeedsConversion) {
auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
Result &= BuildMI(*I.getParent(), I, DL, TII.get(ConvertOpcode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(SignReg)
.constrainAllUses(TII, TRI, RBI);
}
return Result;
}
bool SPIRVInstructionSelector::selectWaveOpInst(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
unsigned Opcode) const {
MachineBasicBlock &BB = *I.getParent();
SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
auto BMI = BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I,
IntTy, TII, !STI.isShader()));
for (unsigned J = 2; J < I.getNumOperands(); J++) {
BMI.addUse(I.getOperand(J).getReg());
}
return BMI.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectWaveActiveCountBits(
Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
SPIRVType *BallotType = GR.getOrCreateSPIRVVectorType(IntTy, 4, I, TII);
Register BallotReg = MRI->createVirtualRegister(GR.getRegClass(BallotType));
bool Result = selectWaveOpInst(BallotReg, BallotType, I,
SPIRV::OpGroupNonUniformBallot);
MachineBasicBlock &BB = *I.getParent();
Result &= BuildMI(BB, I, I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy,
TII, !STI.isShader()))
.addImm(SPIRV::GroupOperation::Reduce)
.addUse(BallotReg)
.constrainAllUses(TII, TRI, RBI);
return Result;
}
bool SPIRVInstructionSelector::selectWaveReduceMax(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
bool IsUnsigned) const {
assert(I.getNumOperands() == 3);
assert(I.getOperand(2).isReg());
MachineBasicBlock &BB = *I.getParent();
Register InputRegister = I.getOperand(2).getReg();
SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
if (!InputType)
report_fatal_error("Input Type could not be determined.");
SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
// Retreive the operation to use based on input type
bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
auto IntegerOpcodeType =
IsUnsigned ? SPIRV::OpGroupNonUniformUMax : SPIRV::OpGroupNonUniformSMax;
auto Opcode = IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntegerOpcodeType;
return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,
!STI.isShader()))
.addImm(SPIRV::GroupOperation::Reduce)
.addUse(I.getOperand(2).getReg())
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectWaveReduceMin(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
bool IsUnsigned) const {
assert(I.getNumOperands() == 3);
assert(I.getOperand(2).isReg());
MachineBasicBlock &BB = *I.getParent();
Register InputRegister = I.getOperand(2).getReg();
SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
if (!InputType)
report_fatal_error("Input Type could not be determined.");
SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
// Retreive the operation to use based on input type
bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
auto IntegerOpcodeType =
IsUnsigned ? SPIRV::OpGroupNonUniformUMin : SPIRV::OpGroupNonUniformSMin;
auto Opcode = IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntegerOpcodeType;
return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,
!STI.isShader()))
.addImm(SPIRV::GroupOperation::Reduce)
.addUse(I.getOperand(2).getReg())
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectWaveReduceSum(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
assert(I.getNumOperands() == 3);
assert(I.getOperand(2).isReg());
MachineBasicBlock &BB = *I.getParent();
Register InputRegister = I.getOperand(2).getReg();
SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
if (!InputType)
report_fatal_error("Input Type could not be determined.");
SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
// Retreive the operation to use based on input type
bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
auto Opcode =
IsFloatTy ? SPIRV::OpGroupNonUniformFAdd : SPIRV::OpGroupNonUniformIAdd;
return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,
!STI.isShader()))
.addImm(SPIRV::GroupOperation::Reduce)
.addUse(I.getOperand(2).getReg());
}
bool SPIRVInstructionSelector::selectBitreverse(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
MachineBasicBlock &BB = *I.getParent();
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpBitReverse))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(1).getReg())
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectFreeze(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
// There is no way to implement `freeze` correctly without support on SPIR-V
// standard side, but we may at least address a simple (static) case when
// undef/poison value presence is obvious. The main benefit of even
// incomplete `freeze` support is preventing of translation from crashing due
// to lack of support on legalization and instruction selection steps.
if (!I.getOperand(0).isReg() || !I.getOperand(1).isReg())
return false;
Register OpReg = I.getOperand(1).getReg();
if (MachineInstr *Def = MRI->getVRegDef(OpReg)) {
if (Def->getOpcode() == TargetOpcode::COPY)
Def = MRI->getVRegDef(Def->getOperand(1).getReg());
Register Reg;
switch (Def->getOpcode()) {
case SPIRV::ASSIGN_TYPE:
if (MachineInstr *AssignToDef =
MRI->getVRegDef(Def->getOperand(1).getReg())) {
if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
Reg = Def->getOperand(2).getReg();
}
break;
case SPIRV::OpUndef:
Reg = Def->getOperand(1).getReg();
break;
}
unsigned DestOpCode;
if (Reg.isValid()) {
DestOpCode = SPIRV::OpConstantNull;
} else {
DestOpCode = TargetOpcode::COPY;
Reg = OpReg;
}
return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(DestOpCode))
.addDef(I.getOperand(0).getReg())
.addUse(Reg)
.constrainAllUses(TII, TRI, RBI);
}
return false;
}
bool SPIRVInstructionSelector::selectBuildVector(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
unsigned N = 0;
if (ResType->getOpcode() == SPIRV::OpTypeVector)
N = GR.getScalarOrVectorComponentCount(ResType);
else if (ResType->getOpcode() == SPIRV::OpTypeArray)
N = getArrayComponentCount(MRI, ResType);
else
report_fatal_error("Cannot select G_BUILD_VECTOR with a non-vector result");
if (I.getNumExplicitOperands() - I.getNumExplicitDefs() != N)
report_fatal_error("G_BUILD_VECTOR and the result type are inconsistent");
// check if we may construct a constant vector
bool IsConst = true;
for (unsigned i = I.getNumExplicitDefs();
i < I.getNumExplicitOperands() && IsConst; ++i)
if (!isConstReg(MRI, I.getOperand(i).getReg()))
IsConst = false;
if (!IsConst && N < 2)
report_fatal_error(
"There must be at least two constituent operands in a vector");
MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(IsConst ? SPIRV::OpConstantComposite
: SPIRV::OpCompositeConstruct))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType));
for (unsigned i = I.getNumExplicitDefs(); i < I.getNumExplicitOperands(); ++i)
MIB.addUse(I.getOperand(i).getReg());
return MIB.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectSplatVector(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
unsigned N = 0;
if (ResType->getOpcode() == SPIRV::OpTypeVector)
N = GR.getScalarOrVectorComponentCount(ResType);
else if (ResType->getOpcode() == SPIRV::OpTypeArray)
N = getArrayComponentCount(MRI, ResType);
else
report_fatal_error("Cannot select G_SPLAT_VECTOR with a non-vector result");
unsigned OpIdx = I.getNumExplicitDefs();
if (!I.getOperand(OpIdx).isReg())
report_fatal_error("Unexpected argument in G_SPLAT_VECTOR");
// check if we may construct a constant vector
Register OpReg = I.getOperand(OpIdx).getReg();
bool IsConst = isConstReg(MRI, OpReg);
if (!IsConst && N < 2)
report_fatal_error(
"There must be at least two constituent operands in a vector");
MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(IsConst ? SPIRV::OpConstantComposite
: SPIRV::OpCompositeConstruct))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType));
for (unsigned i = 0; i < N; ++i)
MIB.addUse(OpReg);
return MIB.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectDiscard(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
unsigned Opcode;
if (STI.canUseExtension(
SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
STI.isAtLeastSPIRVVer(llvm::VersionTuple(1, 6))) {
Opcode = SPIRV::OpDemoteToHelperInvocation;
} else {
Opcode = SPIRV::OpKill;
// OpKill must be the last operation of any basic block.
if (MachineInstr *NextI = I.getNextNode()) {
GR.invalidateMachineInstr(NextI);
NextI->removeFromParent();
}
}
MachineBasicBlock &BB = *I.getParent();
return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectCmp(Register ResVReg,
const SPIRVType *ResType,
unsigned CmpOpc,
MachineInstr &I) const {
Register Cmp0 = I.getOperand(2).getReg();
Register Cmp1 = I.getOperand(3).getReg();
assert(GR.getSPIRVTypeForVReg(Cmp0)->getOpcode() ==
GR.getSPIRVTypeForVReg(Cmp1)->getOpcode() &&
"CMP operands should have the same type");
return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(CmpOpc))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(Cmp0)
.addUse(Cmp1)
.setMIFlags(I.getFlags())
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectICmp(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
auto Pred = I.getOperand(1).getPredicate();
unsigned CmpOpc;
Register CmpOperand = I.getOperand(2).getReg();
if (GR.isScalarOfType(CmpOperand, SPIRV::OpTypePointer))
CmpOpc = getPtrCmpOpcode(Pred);
else if (GR.isScalarOrVectorOfType(CmpOperand, SPIRV::OpTypeBool))
CmpOpc = getBoolCmpOpcode(Pred);
else
CmpOpc = getICmpOpcode(Pred);
return selectCmp(ResVReg, ResType, CmpOpc, I);
}
std::pair<Register, bool>
SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &I,
const SPIRVType *ResType) const {
Type *LLVMTy = IntegerType::get(GR.CurMF->getFunction().getContext(), 32);
const SPIRVType *SpvI32Ty =
ResType ? ResType : GR.getOrCreateSPIRVIntegerType(32, I, TII);
// Find a constant in DT or build a new one.
auto ConstInt = ConstantInt::get(LLVMTy, Val);
Register NewReg = GR.find(ConstInt, GR.CurMF);
bool Result = true;
if (!NewReg.isValid()) {
NewReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
MachineBasicBlock &BB = *I.getParent();
MachineInstr *MI =
Val == 0
? BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
.addDef(NewReg)
.addUse(GR.getSPIRVTypeID(SpvI32Ty))
: BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantI))
.addDef(NewReg)
.addUse(GR.getSPIRVTypeID(SpvI32Ty))
.addImm(APInt(32, Val).getZExtValue());
Result &= constrainSelectedInstRegOperands(*MI, TII, TRI, RBI);
GR.add(ConstInt, MI);
}
return {NewReg, Result};
}
bool SPIRVInstructionSelector::selectFCmp(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
unsigned CmpOp = getFCmpOpcode(I.getOperand(1).getPredicate());
return selectCmp(ResVReg, ResType, CmpOp, I);
}
Register SPIRVInstructionSelector::buildZerosVal(const SPIRVType *ResType,
MachineInstr &I) const {
// OpenCL uses nulls for Zero. In HLSL we don't use null constants.
bool ZeroAsNull = !STI.isShader();
if (ResType->getOpcode() == SPIRV::OpTypeVector)
return GR.getOrCreateConstVector(0UL, I, ResType, TII, ZeroAsNull);
return GR.getOrCreateConstInt(0, I, ResType, TII, ZeroAsNull);
}
Register SPIRVInstructionSelector::buildZerosValF(const SPIRVType *ResType,
MachineInstr &I) const {
// OpenCL uses nulls for Zero. In HLSL we don't use null constants.
bool ZeroAsNull = !STI.isShader();
APFloat VZero = getZeroFP(GR.getTypeForSPIRVType(ResType));
if (ResType->getOpcode() == SPIRV::OpTypeVector)
return GR.getOrCreateConstVector(VZero, I, ResType, TII, ZeroAsNull);
return GR.getOrCreateConstFP(VZero, I, ResType, TII, ZeroAsNull);
}
Register SPIRVInstructionSelector::buildOnesValF(const SPIRVType *ResType,
MachineInstr &I) const {
// OpenCL uses nulls for Zero. In HLSL we don't use null constants.
bool ZeroAsNull = !STI.isShader();
APFloat VOne = getOneFP(GR.getTypeForSPIRVType(ResType));
if (ResType->getOpcode() == SPIRV::OpTypeVector)
return GR.getOrCreateConstVector(VOne, I, ResType, TII, ZeroAsNull);
return GR.getOrCreateConstFP(VOne, I, ResType, TII, ZeroAsNull);
}
Register SPIRVInstructionSelector::buildOnesVal(bool AllOnes,
const SPIRVType *ResType,
MachineInstr &I) const {
unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
APInt One =
AllOnes ? APInt::getAllOnes(BitWidth) : APInt::getOneBitSet(BitWidth, 0);
if (ResType->getOpcode() == SPIRV::OpTypeVector)
return GR.getOrCreateConstVector(One.getZExtValue(), I, ResType, TII);
return GR.getOrCreateConstInt(One.getZExtValue(), I, ResType, TII);
}
bool SPIRVInstructionSelector::selectSelect(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
Register SelectFirstArg = I.getOperand(2).getReg();
Register SelectSecondArg = I.getOperand(3).getReg();
assert(ResType == GR.getSPIRVTypeForVReg(SelectFirstArg) &&
ResType == GR.getSPIRVTypeForVReg(SelectSecondArg));
bool IsFloatTy =
GR.isScalarOrVectorOfType(SelectFirstArg, SPIRV::OpTypeFloat);
bool IsPtrTy =
GR.isScalarOrVectorOfType(SelectFirstArg, SPIRV::OpTypePointer);
bool IsVectorTy = GR.getSPIRVTypeForVReg(SelectFirstArg)->getOpcode() ==
SPIRV::OpTypeVector;
bool IsScalarBool =
GR.isScalarOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool);
unsigned Opcode;
if (IsVectorTy) {
if (IsFloatTy) {
Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
} else if (IsPtrTy) {
Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
} else {
Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
}
} else {
if (IsFloatTy) {
Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
} else if (IsPtrTy) {
Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
} else {
Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
}
}
return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(1).getReg())
.addUse(SelectFirstArg)
.addUse(SelectSecondArg)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectSelectDefaultArgs(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
bool IsSigned) const {
// To extend a bool, we need to use OpSelect between constants.
Register ZeroReg = buildZerosVal(ResType, I);
Register OneReg = buildOnesVal(IsSigned, ResType, I);
bool IsScalarBool =
GR.isScalarOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool);
unsigned Opcode =
IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(1).getReg())
.addUse(OneReg)
.addUse(ZeroReg)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectIToF(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I, bool IsSigned,
unsigned Opcode) const {
Register SrcReg = I.getOperand(1).getReg();
// We can convert bool value directly to float type without OpConvert*ToF,
// however the translator generates OpSelect+OpConvert*ToF, so we do the same.
if (GR.isScalarOrVectorOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool)) {
unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
SPIRVType *TmpType = GR.getOrCreateSPIRVIntegerType(BitWidth, I, TII);
if (ResType->getOpcode() == SPIRV::OpTypeVector) {
const unsigned NumElts = ResType->getOperand(2).getImm();
TmpType = GR.getOrCreateSPIRVVectorType(TmpType, NumElts, I, TII);
}
SrcReg = createVirtualRegister(TmpType, &GR, MRI, MRI->getMF());
selectSelectDefaultArgs(SrcReg, TmpType, I, false);
}
return selectOpWithSrcs(ResVReg, ResType, I, {SrcReg}, Opcode);
}
bool SPIRVInstructionSelector::selectExt(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I, bool IsSigned) const {
Register SrcReg = I.getOperand(1).getReg();
if (GR.isScalarOrVectorOfType(SrcReg, SPIRV::OpTypeBool))
return selectSelectDefaultArgs(ResVReg, ResType, I, IsSigned);
SPIRVType *SrcType = GR.getSPIRVTypeForVReg(SrcReg);
if (SrcType == ResType)
return BuildCOPY(ResVReg, SrcReg, I);
unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
return selectUnOp(ResVReg, ResType, I, Opcode);
}
bool SPIRVInstructionSelector::selectSUCmp(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
bool IsSigned) const {
MachineIRBuilder MIRBuilder(I);
MachineRegisterInfo *MRI = MIRBuilder.getMRI();
MachineBasicBlock &BB = *I.getParent();
// Ensure we have bool.
SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
unsigned N = GR.getScalarOrVectorComponentCount(ResType);
if (N > 1)
BoolType = GR.getOrCreateSPIRVVectorType(BoolType, N, I, TII);
Register BoolTypeReg = GR.getSPIRVTypeID(BoolType);
// Build less-than-equal and less-than.
// TODO: replace with one-liner createVirtualRegister() from
// llvm/lib/Target/SPIRV/SPIRVUtils.cpp when PR #116609 is merged.
Register IsLessEqReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
MRI->setType(IsLessEqReg, LLT::scalar(64));
GR.assignSPIRVTypeToVReg(ResType, IsLessEqReg, MIRBuilder.getMF());
bool Result = BuildMI(BB, I, I.getDebugLoc(),
TII.get(IsSigned ? SPIRV::OpSLessThanEqual
: SPIRV::OpULessThanEqual))
.addDef(IsLessEqReg)
.addUse(BoolTypeReg)
.addUse(I.getOperand(1).getReg())
.addUse(I.getOperand(2).getReg())
.constrainAllUses(TII, TRI, RBI);
Register IsLessReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
MRI->setType(IsLessReg, LLT::scalar(64));
GR.assignSPIRVTypeToVReg(ResType, IsLessReg, MIRBuilder.getMF());
Result &= BuildMI(BB, I, I.getDebugLoc(),
TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
.addDef(IsLessReg)
.addUse(BoolTypeReg)
.addUse(I.getOperand(1).getReg())
.addUse(I.getOperand(2).getReg())
.constrainAllUses(TII, TRI, RBI);
// Build selects.
Register ResTypeReg = GR.getSPIRVTypeID(ResType);
Register NegOneOrZeroReg =
MRI->createVirtualRegister(GR.getRegClass(ResType));
MRI->setType(NegOneOrZeroReg, LLT::scalar(64));
GR.assignSPIRVTypeToVReg(ResType, NegOneOrZeroReg, MIRBuilder.getMF());
unsigned SelectOpcode =
N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SelectOpcode))
.addDef(NegOneOrZeroReg)
.addUse(ResTypeReg)
.addUse(IsLessReg)
.addUse(buildOnesVal(true, ResType, I)) // -1
.addUse(buildZerosVal(ResType, I))
.constrainAllUses(TII, TRI, RBI);
return Result & BuildMI(BB, I, I.getDebugLoc(), TII.get(SelectOpcode))
.addDef(ResVReg)
.addUse(ResTypeReg)
.addUse(IsLessEqReg)
.addUse(NegOneOrZeroReg) // -1 or 0
.addUse(buildOnesVal(false, ResType, I))
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectIntToBool(Register IntReg,
Register ResVReg,
MachineInstr &I,
const SPIRVType *IntTy,
const SPIRVType *BoolTy) const {
// To truncate to a bool, we use OpBitwiseAnd 1 and OpINotEqual to zero.
Register BitIntReg = createVirtualRegister(IntTy, &GR, MRI, MRI->getMF());
bool IsVectorTy = IntTy->getOpcode() == SPIRV::OpTypeVector;
unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
Register Zero = buildZerosVal(IntTy, I);
Register One = buildOnesVal(false, IntTy, I);
MachineBasicBlock &BB = *I.getParent();
bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
.addDef(BitIntReg)
.addUse(GR.getSPIRVTypeID(IntTy))
.addUse(IntReg)
.addUse(One)
.constrainAllUses(TII, TRI, RBI);
return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpINotEqual))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(BoolTy))
.addUse(BitIntReg)
.addUse(Zero)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectTrunc(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
Register IntReg = I.getOperand(1).getReg();
const SPIRVType *ArgType = GR.getSPIRVTypeForVReg(IntReg);
if (GR.isScalarOrVectorOfType(ResVReg, SPIRV::OpTypeBool))
return selectIntToBool(IntReg, ResVReg, I, ArgType, ResType);
if (ArgType == ResType)
return BuildCOPY(ResVReg, IntReg, I);
bool IsSigned = GR.isScalarOrVectorSigned(ResType);
unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
return selectUnOp(ResVReg, ResType, I, Opcode);
}
bool SPIRVInstructionSelector::selectConst(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
unsigned Opcode = I.getOpcode();
unsigned TpOpcode = ResType->getOpcode();
Register Reg;
if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
assert(Opcode == TargetOpcode::G_CONSTANT &&
I.getOperand(1).getCImm()->isZero());
MachineBasicBlock &DepMBB = I.getMF()->front();
MachineIRBuilder MIRBuilder(DepMBB, DepMBB.getFirstNonPHI());
Reg = GR.getOrCreateConstNullPtr(MIRBuilder, ResType);
} else if (Opcode == TargetOpcode::G_FCONSTANT) {
Reg = GR.getOrCreateConstFP(I.getOperand(1).getFPImm()->getValue(), I,
ResType, TII, !STI.isShader());
} else {
Reg = GR.getOrCreateConstInt(I.getOperand(1).getCImm()->getZExtValue(), I,
ResType, TII, !STI.isShader());
}
return Reg == ResVReg ? true : BuildCOPY(ResVReg, Reg, I);
}
bool SPIRVInstructionSelector::selectOpUndef(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectInsertVal(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
MachineBasicBlock &BB = *I.getParent();
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeInsert))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
// object to insert
.addUse(I.getOperand(3).getReg())
// composite to insert into
.addUse(I.getOperand(2).getReg());
for (unsigned i = 4; i < I.getNumOperands(); i++)
MIB.addImm(foldImm(I.getOperand(i), MRI));
return MIB.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectExtractVal(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
MachineBasicBlock &BB = *I.getParent();
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(2).getReg());
for (unsigned i = 3; i < I.getNumOperands(); i++)
MIB.addImm(foldImm(I.getOperand(i), MRI));
return MIB.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectInsertElt(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
if (getImm(I.getOperand(4), MRI))
return selectInsertVal(ResVReg, ResType, I);
MachineBasicBlock &BB = *I.getParent();
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorInsertDynamic))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(2).getReg())
.addUse(I.getOperand(3).getReg())
.addUse(I.getOperand(4).getReg())
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectExtractElt(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
if (getImm(I.getOperand(3), MRI))
return selectExtractVal(ResVReg, ResType, I);
MachineBasicBlock &BB = *I.getParent();
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorExtractDynamic))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(2).getReg())
.addUse(I.getOperand(3).getReg())
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectGEP(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
const bool IsGEPInBounds = I.getOperand(2).getImm();
// OpAccessChain could be used for OpenCL, but the SPIRV-LLVM Translator only
// relies on PtrAccessChain, so we'll try not to deviate. For Vulkan however,
// we have to use Op[InBounds]AccessChain.
const unsigned Opcode = STI.isLogicalSPIRV()
? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
: SPIRV::OpAccessChain)
: (IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
: SPIRV::OpPtrAccessChain);
auto Res = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
// Object to get a pointer to.
.addUse(I.getOperand(3).getReg());
assert(
(Opcode == SPIRV::OpPtrAccessChain ||
Opcode == SPIRV::OpInBoundsPtrAccessChain ||
(getImm(I.getOperand(4), MRI) && foldImm(I.getOperand(4), MRI) == 0)) &&
"Cannot translate GEP to OpAccessChain. First index must be 0.");
// Adding indices.
const unsigned StartingIndex =
(Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
? 5
: 4;
for (unsigned i = StartingIndex; i < I.getNumExplicitOperands(); ++i)
Res.addUse(I.getOperand(i).getReg());
return Res.constrainAllUses(TII, TRI, RBI);
}
// Maybe wrap a value into OpSpecConstantOp
bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
MachineInstr &I, SmallVector<Register> &CompositeArgs) const {
bool Result = true;
unsigned Lim = I.getNumExplicitOperands();
for (unsigned i = I.getNumExplicitDefs() + 1; i < Lim; ++i) {
Register OpReg = I.getOperand(i).getReg();
MachineInstr *OpDefine = MRI->getVRegDef(OpReg);
SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
SmallPtrSet<SPIRVType *, 4> Visited;
if (!OpDefine || !OpType || isConstReg(MRI, OpDefine, Visited) ||
OpDefine->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
OpDefine->getOpcode() == TargetOpcode::G_INTTOPTR ||
GR.isAggregateType(OpType)) {
// The case of G_ADDRSPACE_CAST inside spv_const_composite() is processed
// by selectAddrSpaceCast(), and G_INTTOPTR is processed by selectUnOp()
CompositeArgs.push_back(OpReg);
continue;
}
MachineFunction *MF = I.getMF();
Register WrapReg = GR.find(OpDefine, MF);
if (WrapReg.isValid()) {
CompositeArgs.push_back(WrapReg);
continue;
}
// Create a new register for the wrapper
WrapReg = MRI->createVirtualRegister(GR.getRegClass(OpType));
CompositeArgs.push_back(WrapReg);
// Decorate the wrapper register and generate a new instruction
MRI->setType(WrapReg, LLT::pointer(0, 64));
GR.assignSPIRVTypeToVReg(OpType, WrapReg, *MF);
auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
.addDef(WrapReg)
.addUse(GR.getSPIRVTypeID(OpType))
.addImm(static_cast<uint32_t>(SPIRV::Opcode::Bitcast))
.addUse(OpReg);
GR.add(OpDefine, MIB);
Result = MIB.constrainAllUses(TII, TRI, RBI);
if (!Result)
break;
}
return Result;
}
bool SPIRVInstructionSelector::selectDerivativeInst(
Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
const unsigned DPdOpCode) const {
// TODO: This should check specifically for Fragment Execution Model, but STI
// doesn't provide that information yet. See #167562
errorIfInstrOutsideShader(I);
// If the arg/result types are half then we need to wrap the instr in
// conversions to float
// This case occurs because a half arg/result is legal in HLSL but not spirv.
Register SrcReg = I.getOperand(2).getReg();
SPIRVType *SrcType = GR.getSPIRVTypeForVReg(SrcReg);
unsigned BitWidth = std::min(GR.getScalarOrVectorBitWidth(SrcType),
GR.getScalarOrVectorBitWidth(ResType));
if (BitWidth == 32)
return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(DPdOpCode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(2).getReg());
MachineIRBuilder MIRBuilder(I);
unsigned componentCount = GR.getScalarOrVectorComponentCount(SrcType);
SPIRVType *F32ConvertTy = GR.getOrCreateSPIRVFloatType(32, I, TII);
if (componentCount != 1)
F32ConvertTy = GR.getOrCreateSPIRVVectorType(F32ConvertTy, componentCount,
MIRBuilder, false);
const TargetRegisterClass *RegClass = GR.getRegClass(SrcType);
Register ConvertToVReg = MRI->createVirtualRegister(RegClass);
Register DpdOpVReg = MRI->createVirtualRegister(RegClass);
bool Result =
BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpFConvert))
.addDef(ConvertToVReg)
.addUse(GR.getSPIRVTypeID(F32ConvertTy))
.addUse(SrcReg)
.constrainAllUses(TII, TRI, RBI);
Result &= BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(DPdOpCode))
.addDef(DpdOpVReg)
.addUse(GR.getSPIRVTypeID(F32ConvertTy))
.addUse(ConvertToVReg)
.constrainAllUses(TII, TRI, RBI);
Result &=
BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpFConvert))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(DpdOpVReg)
.constrainAllUses(TII, TRI, RBI);
return Result;
}
bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
MachineBasicBlock &BB = *I.getParent();
Intrinsic::ID IID = cast<GIntrinsic>(I).getIntrinsicID();
switch (IID) {
case Intrinsic::spv_load:
return selectLoad(ResVReg, ResType, I);
case Intrinsic::spv_store:
return selectStore(I);
case Intrinsic::spv_extractv:
return selectExtractVal(ResVReg, ResType, I);
case Intrinsic::spv_insertv:
return selectInsertVal(ResVReg, ResType, I);
case Intrinsic::spv_extractelt:
return selectExtractElt(ResVReg, ResType, I);
case Intrinsic::spv_insertelt:
return selectInsertElt(ResVReg, ResType, I);
case Intrinsic::spv_gep:
return selectGEP(ResVReg, ResType, I);
case Intrinsic::spv_bitcast: {
Register OpReg = I.getOperand(2).getReg();
SPIRVType *OpType =
OpReg.isValid() ? GR.getSPIRVTypeForVReg(OpReg) : nullptr;
if (!GR.isBitcastCompatible(ResType, OpType))
report_fatal_error("incompatible result and operand types in a bitcast");
return selectOpWithSrcs(ResVReg, ResType, I, {OpReg}, SPIRV::OpBitcast);
}
case Intrinsic::spv_unref_global:
case Intrinsic::spv_init_global: {
MachineInstr *MI = MRI->getVRegDef(I.getOperand(1).getReg());
MachineInstr *Init = I.getNumExplicitOperands() > 2
? MRI->getVRegDef(I.getOperand(2).getReg())
: nullptr;
assert(MI);
Register GVarVReg = MI->getOperand(0).getReg();
bool Res = selectGlobalValue(GVarVReg, *MI, Init);
// We violate SSA form by inserting OpVariable and still having a gMIR
// instruction %vreg = G_GLOBAL_VALUE @gvar. We need to fix this by erasing
// the duplicated definition.
if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
GR.invalidateMachineInstr(MI);
MI->removeFromParent();
}
return Res;
}
case Intrinsic::spv_undef: {
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType));
return MIB.constrainAllUses(TII, TRI, RBI);
}
case Intrinsic::spv_const_composite: {
// If no values are attached, the composite is null constant.
bool IsNull = I.getNumExplicitDefs() + 1 == I.getNumExplicitOperands();
SmallVector<Register> CompositeArgs;
MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
// skip type MD node we already used when generated assign.type for this
if (!IsNull) {
if (!wrapIntoSpecConstantOp(I, CompositeArgs))
return false;
MachineIRBuilder MIR(I);
SmallVector<MachineInstr *, 4> Instructions = createContinuedInstructions(
MIR, SPIRV::OpConstantComposite, 3,
SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,
GR.getSPIRVTypeID(ResType));
for (auto *Instr : Instructions) {
Instr->setDebugLoc(I.getDebugLoc());
if (!constrainSelectedInstRegOperands(*Instr, TII, TRI, RBI))
return false;
}
return true;
} else {
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType));
return MIB.constrainAllUses(TII, TRI, RBI);
}
}
case Intrinsic::spv_assign_name: {
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpName));
MIB.addUse(I.getOperand(I.getNumExplicitDefs() + 1).getReg());
for (unsigned i = I.getNumExplicitDefs() + 2;
i < I.getNumExplicitOperands(); ++i) {
MIB.addImm(I.getOperand(i).getImm());
}
return MIB.constrainAllUses(TII, TRI, RBI);
}
case Intrinsic::spv_switch: {
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSwitch));
for (unsigned i = 1; i < I.getNumExplicitOperands(); ++i) {
if (I.getOperand(i).isReg())
MIB.addReg(I.getOperand(i).getReg());
else if (I.getOperand(i).isCImm())
addNumImm(I.getOperand(i).getCImm()->getValue(), MIB);
else if (I.getOperand(i).isMBB())
MIB.addMBB(I.getOperand(i).getMBB());
else
llvm_unreachable("Unexpected OpSwitch operand");
}
return MIB.constrainAllUses(TII, TRI, RBI);
}
case Intrinsic::spv_loop_merge: {
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpLoopMerge));
for (unsigned i = 1; i < I.getNumExplicitOperands(); ++i) {
if (I.getOperand(i).isMBB())
MIB.addMBB(I.getOperand(i).getMBB());
else
MIB.addImm(foldImm(I.getOperand(i), MRI));
}
return MIB.constrainAllUses(TII, TRI, RBI);
}
case Intrinsic::spv_selection_merge: {
auto MIB =
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSelectionMerge));
assert(I.getOperand(1).isMBB() &&
"operand 1 to spv_selection_merge must be a basic block");
MIB.addMBB(I.getOperand(1).getMBB());
MIB.addImm(getSelectionOperandForImm(I.getOperand(2).getImm()));
return MIB.constrainAllUses(TII, TRI, RBI);
}
case Intrinsic::spv_cmpxchg:
return selectAtomicCmpXchg(ResVReg, ResType, I);
case Intrinsic::spv_unreachable:
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUnreachable))
.constrainAllUses(TII, TRI, RBI);
case Intrinsic::spv_alloca:
return selectFrameIndex(ResVReg, ResType, I);
case Intrinsic::spv_alloca_array:
return selectAllocaArray(ResVReg, ResType, I);
case Intrinsic::spv_assume:
if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpAssumeTrueKHR))
.addUse(I.getOperand(1).getReg())
.constrainAllUses(TII, TRI, RBI);
break;
case Intrinsic::spv_expect:
if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExpectKHR))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(2).getReg())
.addUse(I.getOperand(3).getReg())
.constrainAllUses(TII, TRI, RBI);
break;
case Intrinsic::arithmetic_fence:
if (STI.canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence))
return BuildMI(BB, I, I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(2).getReg())
.constrainAllUses(TII, TRI, RBI);
else
return BuildCOPY(ResVReg, I.getOperand(2).getReg(), I);
break;
case Intrinsic::spv_thread_id:
// The HLSL SV_DispatchThreadID semantic is lowered to llvm.spv.thread.id
// intrinsic in LLVM IR for SPIR-V backend.
//
// In SPIR-V backend, llvm.spv.thread.id is now correctly translated to a
// `GlobalInvocationId` builtin variable
return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
ResType, I);
case Intrinsic::spv_thread_id_in_group:
// The HLSL SV_GroupThreadId semantic is lowered to
// llvm.spv.thread.id.in.group intrinsic in LLVM IR for SPIR-V backend.
//
// In SPIR-V backend, llvm.spv.thread.id.in.group is now correctly
// translated to a `LocalInvocationId` builtin variable
return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
ResType, I);
case Intrinsic::spv_group_id:
// The HLSL SV_GroupId semantic is lowered to
// llvm.spv.group.id intrinsic in LLVM IR for SPIR-V backend.
//
// In SPIR-V backend, llvm.spv.group.id is now translated to a `WorkgroupId`
// builtin variable
return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
I);
case Intrinsic::spv_flattened_thread_id_in_group:
// The HLSL SV_GroupIndex semantic is lowered to
// llvm.spv.flattened.thread.id.in.group() intrinsic in LLVM IR for SPIR-V
// backend.
//
// In SPIR-V backend, llvm.spv.flattened.thread.id.in.group is translated to
// a `LocalInvocationIndex` builtin variable
return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
ResType, I);
case Intrinsic::spv_workgroup_size:
return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
ResType, I);
case Intrinsic::spv_global_size:
return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
I);
case Intrinsic::spv_global_offset:
return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
ResType, I);
case Intrinsic::spv_num_workgroups:
return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
ResType, I);
case Intrinsic::spv_subgroup_size:
return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
I);
case Intrinsic::spv_num_subgroups:
return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
I);
case Intrinsic::spv_subgroup_id:
return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType, I);
case Intrinsic::spv_subgroup_local_invocation_id:
return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
ResVReg, ResType, I);
case Intrinsic::spv_subgroup_max_size:
return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
I);
case Intrinsic::spv_fdot:
return selectFloatDot(ResVReg, ResType, I);
case Intrinsic::spv_udot:
case Intrinsic::spv_sdot:
if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
return selectIntegerDot(ResVReg, ResType, I,
/*Signed=*/IID == Intrinsic::spv_sdot);
return selectIntegerDotExpansion(ResVReg, ResType, I);
case Intrinsic::spv_dot4add_i8packed:
if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
return selectDot4AddPacked<true>(ResVReg, ResType, I);
return selectDot4AddPackedExpansion<true>(ResVReg, ResType, I);
case Intrinsic::spv_dot4add_u8packed:
if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
return selectDot4AddPacked<false>(ResVReg, ResType, I);
return selectDot4AddPackedExpansion<false>(ResVReg, ResType, I);
case Intrinsic::spv_all:
return selectAll(ResVReg, ResType, I);
case Intrinsic::spv_any:
return selectAny(ResVReg, ResType, I);
case Intrinsic::spv_cross:
return selectExtInst(ResVReg, ResType, I, CL::cross, GL::Cross);
case Intrinsic::spv_distance:
return selectExtInst(ResVReg, ResType, I, CL::distance, GL::Distance);
case Intrinsic::spv_lerp:
return selectExtInst(ResVReg, ResType, I, CL::mix, GL::FMix);
case Intrinsic::spv_length:
return selectExtInst(ResVReg, ResType, I, CL::length, GL::Length);
case Intrinsic::spv_degrees:
return selectExtInst(ResVReg, ResType, I, CL::degrees, GL::Degrees);
case Intrinsic::spv_faceforward:
return selectExtInst(ResVReg, ResType, I, GL::FaceForward);
case Intrinsic::spv_frac:
return selectExtInst(ResVReg, ResType, I, CL::fract, GL::Fract);
case Intrinsic::spv_isinf:
return selectOpIsInf(ResVReg, ResType, I);
case Intrinsic::spv_isnan:
return selectOpIsNan(ResVReg, ResType, I);
case Intrinsic::spv_normalize:
return selectExtInst(ResVReg, ResType, I, CL::normalize, GL::Normalize);
case Intrinsic::spv_refract:
return selectExtInst(ResVReg, ResType, I, GL::Refract);
case Intrinsic::spv_reflect:
return selectExtInst(ResVReg, ResType, I, GL::Reflect);
case Intrinsic::spv_rsqrt:
return selectExtInst(ResVReg, ResType, I, CL::rsqrt, GL::InverseSqrt);
case Intrinsic::spv_sign:
return selectSign(ResVReg, ResType, I);
case Intrinsic::spv_smoothstep:
return selectExtInst(ResVReg, ResType, I, CL::smoothstep, GL::SmoothStep);
case Intrinsic::spv_firstbituhigh: // There is no CL equivalent of FindUMsb
return selectFirstBitHigh(ResVReg, ResType, I, /*IsSigned=*/false);
case Intrinsic::spv_firstbitshigh: // There is no CL equivalent of FindSMsb
return selectFirstBitHigh(ResVReg, ResType, I, /*IsSigned=*/true);
case Intrinsic::spv_firstbitlow: // There is no CL equivlent of FindILsb
return selectFirstBitLow(ResVReg, ResType, I);
case Intrinsic::spv_group_memory_barrier_with_group_sync: {
bool Result = true;
auto MemSemConstant =
buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent, I);
Register MemSemReg = MemSemConstant.first;
Result &= MemSemConstant.second;
auto ScopeConstant = buildI32Constant(SPIRV::Scope::Workgroup, I);
Register ScopeReg = ScopeConstant.first;
Result &= ScopeConstant.second;
MachineBasicBlock &BB = *I.getParent();
return Result &&
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpControlBarrier))
.addUse(ScopeReg)
.addUse(ScopeReg)
.addUse(MemSemReg)
.constrainAllUses(TII, TRI, RBI);
}
case Intrinsic::spv_generic_cast_to_ptr_explicit: {
Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 1).getReg();
SPIRV::StorageClass::StorageClass ResSC =
GR.getPointerStorageClass(ResType);
if (!isGenericCastablePtr(ResSC))
report_fatal_error("The target storage class is not castable from the "
"Generic storage class");
return BuildMI(BB, I, I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(PtrReg)
.addImm(ResSC)
.constrainAllUses(TII, TRI, RBI);
}
case Intrinsic::spv_lifetime_start:
case Intrinsic::spv_lifetime_end: {
unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
: SPIRV::OpLifetimeStop;
int64_t Size = I.getOperand(I.getNumExplicitDefs() + 1).getImm();
Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 2).getReg();
if (Size == -1)
Size = 0;
return BuildMI(BB, I, I.getDebugLoc(), TII.get(Op))
.addUse(PtrReg)
.addImm(Size)
.constrainAllUses(TII, TRI, RBI);
}
case Intrinsic::spv_saturate:
return selectSaturate(ResVReg, ResType, I);
case Intrinsic::spv_nclamp:
return selectExtInst(ResVReg, ResType, I, CL::fclamp, GL::NClamp);
case Intrinsic::spv_uclamp:
return selectExtInst(ResVReg, ResType, I, CL::u_clamp, GL::UClamp);
case Intrinsic::spv_sclamp:
return selectExtInst(ResVReg, ResType, I, CL::s_clamp, GL::SClamp);
case Intrinsic::spv_wave_active_countbits:
return selectWaveActiveCountBits(ResVReg, ResType, I);
case Intrinsic::spv_wave_all:
return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformAll);
case Intrinsic::spv_wave_any:
return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformAny);
case Intrinsic::spv_wave_is_first_lane:
return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformElect);
case Intrinsic::spv_wave_reduce_umax:
return selectWaveReduceMax(ResVReg, ResType, I, /*IsUnsigned*/ true);
case Intrinsic::spv_wave_reduce_max:
return selectWaveReduceMax(ResVReg, ResType, I, /*IsUnsigned*/ false);
case Intrinsic::spv_wave_reduce_umin:
return selectWaveReduceMin(ResVReg, ResType, I, /*IsUnsigned*/ true);
case Intrinsic::spv_wave_reduce_min:
return selectWaveReduceMin(ResVReg, ResType, I, /*IsUnsigned*/ false);
case Intrinsic::spv_wave_reduce_sum:
return selectWaveReduceSum(ResVReg, ResType, I);
case Intrinsic::spv_wave_readlane:
return selectWaveOpInst(ResVReg, ResType, I,
SPIRV::OpGroupNonUniformShuffle);
case Intrinsic::spv_step:
return selectExtInst(ResVReg, ResType, I, CL::step, GL::Step);
case Intrinsic::spv_radians:
return selectExtInst(ResVReg, ResType, I, CL::radians, GL::Radians);
// Discard intrinsics which we do not expect to actually represent code after
// lowering or intrinsics which are not implemented but should not crash when
// found in a customer's LLVM IR input.
case Intrinsic::instrprof_increment:
case Intrinsic::instrprof_increment_step:
case Intrinsic::instrprof_value_profile:
break;
// Discard internal intrinsics.
case Intrinsic::spv_value_md:
break;
case Intrinsic::spv_resource_handlefrombinding: {
return selectHandleFromBinding(ResVReg, ResType, I);
}
case Intrinsic::spv_resource_counterhandlefrombinding:
return selectCounterHandleFromBinding(ResVReg, ResType, I);
case Intrinsic::spv_resource_updatecounter:
return selectUpdateCounter(ResVReg, ResType, I);
case Intrinsic::spv_resource_store_typedbuffer: {
return selectImageWriteIntrinsic(I);
}
case Intrinsic::spv_resource_load_typedbuffer: {
return selectReadImageIntrinsic(ResVReg, ResType, I);
}
case Intrinsic::spv_resource_getpointer: {
return selectResourceGetPointer(ResVReg, ResType, I);
}
case Intrinsic::spv_pushconstant_getpointer: {
return selectPushConstantGetPointer(ResVReg, ResType, I);
}
case Intrinsic::spv_discard: {
return selectDiscard(ResVReg, ResType, I);
}
case Intrinsic::spv_resource_nonuniformindex: {
return selectResourceNonUniformIndex(ResVReg, ResType, I);
}
case Intrinsic::spv_unpackhalf2x16: {
return selectExtInst(ResVReg, ResType, I, GL::UnpackHalf2x16);
}
case Intrinsic::spv_ddx:
return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdx);
case Intrinsic::spv_ddy:
return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdy);
case Intrinsic::spv_ddx_coarse:
return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdxCoarse);
case Intrinsic::spv_ddy_coarse:
return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdyCoarse);
case Intrinsic::spv_ddx_fine:
return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdxFine);
case Intrinsic::spv_ddy_fine:
return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdyFine);
case Intrinsic::spv_fwidth:
return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpFwidth);
default: {
std::string DiagMsg;
raw_string_ostream OS(DiagMsg);
I.print(OS);
DiagMsg = "Intrinsic selection not implemented: " + DiagMsg;
report_fatal_error(DiagMsg.c_str(), false);
}
}
return true;
}
bool SPIRVInstructionSelector::selectHandleFromBinding(Register &ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
// The images need to be loaded in the same basic block as their use. We defer
// loading the image to the intrinsic that uses it.
if (ResType->getOpcode() == SPIRV::OpTypeImage)
return true;
return loadHandleBeforePosition(ResVReg, GR.getSPIRVTypeForVReg(ResVReg),
*cast<GIntrinsic>(&I), I);
}
bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
auto &Intr = cast<GIntrinsic>(I);
assert(Intr.getIntrinsicID() ==
Intrinsic::spv_resource_counterhandlefrombinding);
// Extract information from the intrinsic call.
Register MainHandleReg = Intr.getOperand(2).getReg();
auto *MainHandleDef = cast<GIntrinsic>(getVRegDef(*MRI, MainHandleReg));
assert(MainHandleDef->getIntrinsicID() ==
Intrinsic::spv_resource_handlefrombinding);
uint32_t Set = getIConstVal(Intr.getOperand(4).getReg(), MRI);
uint32_t Binding = getIConstVal(Intr.getOperand(3).getReg(), MRI);
uint32_t ArraySize = getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
Register IndexReg = MainHandleDef->getOperand(5).getReg();
std::string CounterName =
getStringValueFromReg(MainHandleDef->getOperand(6).getReg(), *MRI) +
".counter";
// Create the counter variable.
MachineIRBuilder MIRBuilder(I);
Register CounterVarReg = buildPointerToResource(
GR.getPointeeType(ResType), GR.getPointerStorageClass(ResType), Set,
Binding, ArraySize, IndexReg, CounterName, MIRBuilder);
return BuildCOPY(ResVReg, CounterVarReg, I);
}
bool SPIRVInstructionSelector::selectUpdateCounter(Register &ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
auto &Intr = cast<GIntrinsic>(I);
assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
Register CounterHandleReg = Intr.getOperand(2).getReg();
Register IncrReg = Intr.getOperand(3).getReg();
// The counter handle is a pointer to the counter variable (which is a struct
// containing an i32). We need to get a pointer to that i32 member to do the
// atomic operation.
#ifndef NDEBUG
SPIRVType *CounterVarType = GR.getSPIRVTypeForVReg(CounterHandleReg);
SPIRVType *CounterVarPointeeType = GR.getPointeeType(CounterVarType);
assert(CounterVarPointeeType &&
CounterVarPointeeType->getOpcode() == SPIRV::OpTypeStruct &&
"Counter variable must be a struct");
assert(GR.getPointerStorageClass(CounterVarType) ==
SPIRV::StorageClass::StorageBuffer &&
"Counter variable must be in the storage buffer storage class");
assert(CounterVarPointeeType->getNumOperands() == 2 &&
"Counter variable must have exactly 1 member in the struct");
const SPIRVType *MemberType =
GR.getSPIRVTypeForVReg(CounterVarPointeeType->getOperand(1).getReg());
assert(MemberType->getOpcode() == SPIRV::OpTypeInt &&
"Counter variable struct must have a single i32 member");
#endif
// The struct has a single i32 member.
MachineIRBuilder MIRBuilder(I);
const Type *LLVMIntType =
Type::getInt32Ty(I.getMF()->getFunction().getContext());
SPIRVType *IntPtrType = GR.getOrCreateSPIRVPointerType(
LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
auto Zero = buildI32Constant(0, I);
if (!Zero.second)
return false;
Register PtrToCounter =
MRI->createVirtualRegister(GR.getRegClass(IntPtrType));
if (!BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
.addDef(PtrToCounter)
.addUse(GR.getSPIRVTypeID(IntPtrType))
.addUse(CounterHandleReg)
.addUse(Zero.first)
.constrainAllUses(TII, TRI, RBI)) {
return false;
}
// For UAV/SSBO counters, the scope is Device. The counter variable is not
// used as a flag. So the memory semantics can be None.
auto Scope = buildI32Constant(SPIRV::Scope::Device, I);
if (!Scope.second)
return false;
auto Semantics = buildI32Constant(SPIRV::MemorySemantics::None, I);
if (!Semantics.second)
return false;
int64_t IncrVal = getIConstValSext(IncrReg, MRI);
auto Incr = buildI32Constant(static_cast<uint32_t>(IncrVal), I);
if (!Incr.second)
return false;
Register AtomicRes = MRI->createVirtualRegister(GR.getRegClass(ResType));
if (!BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpAtomicIAdd))
.addDef(AtomicRes)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(PtrToCounter)
.addUse(Scope.first)
.addUse(Semantics.first)
.addUse(Incr.first)
.constrainAllUses(TII, TRI, RBI)) {
return false;
}
if (IncrVal >= 0) {
return BuildCOPY(ResVReg, AtomicRes, I);
}
// In HLSL, IncrementCounter returns the value *before* the increment, while
// DecrementCounter returns the value *after* the decrement. Both are lowered
// to the same atomic intrinsic which returns the value *before* the
// operation. So for decrements (negative IncrVal), we must subtract the
// increment value from the result to get the post-decrement value.
return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(AtomicRes)
.addUse(Incr.first)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
// If the load of the image is in a different basic block, then
// this will generate invalid code. A proper solution is to move
// the OpLoad from selectHandleFromBinding here. However, to do
// that we will need to change the return type of the intrinsic.
// We will do that when we can, but for now trying to move forward with other
// issues.
Register ImageReg = I.getOperand(2).getReg();
auto *ImageDef = cast<GIntrinsic>(getVRegDef(*MRI, ImageReg));
Register NewImageReg = MRI->createVirtualRegister(MRI->getRegClass(ImageReg));
if (!loadHandleBeforePosition(NewImageReg, GR.getSPIRVTypeForVReg(ImageReg),
*ImageDef, I)) {
return false;
}
Register IdxReg = I.getOperand(3).getReg();
DebugLoc Loc = I.getDebugLoc();
MachineInstr &Pos = I;
return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
Pos);
}
bool SPIRVInstructionSelector::generateImageReadOrFetch(
Register &ResVReg, const SPIRVType *ResType, Register ImageReg,
Register IdxReg, DebugLoc Loc, MachineInstr &Pos) const {
SPIRVType *ImageType = GR.getSPIRVTypeForVReg(ImageReg);
assert(ImageType && ImageType->getOpcode() == SPIRV::OpTypeImage &&
"ImageReg is not an image type.");
bool IsSignedInteger =
sampledTypeIsSignedInteger(GR.getTypeForSPIRVType(ImageType));
// Check if the "sampled" operand of the image type is 1.
// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpImageFetch
auto SampledOp = ImageType->getOperand(6);
bool IsFetch = (SampledOp.getImm() == 1);
uint64_t ResultSize = GR.getScalarOrVectorComponentCount(ResType);
if (ResultSize == 4) {
auto BMI =
BuildMI(*Pos.getParent(), Pos, Loc,
TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(ImageReg)
.addUse(IdxReg);
if (IsSignedInteger)
BMI.addImm(0x1000); // SignExtend
return BMI.constrainAllUses(TII, TRI, RBI);
}
SPIRVType *ReadType = widenTypeToVec4(ResType, Pos);
Register ReadReg = MRI->createVirtualRegister(GR.getRegClass(ReadType));
auto BMI =
BuildMI(*Pos.getParent(), Pos, Loc,
TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
.addDef(ReadReg)
.addUse(GR.getSPIRVTypeID(ReadType))
.addUse(ImageReg)
.addUse(IdxReg);
if (IsSignedInteger)
BMI.addImm(0x1000); // SignExtend
bool Succeed = BMI.constrainAllUses(TII, TRI, RBI);
if (!Succeed)
return false;
if (ResultSize == 1) {
return BuildMI(*Pos.getParent(), Pos, Loc,
TII.get(SPIRV::OpCompositeExtract))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(ReadReg)
.addImm(0)
.constrainAllUses(TII, TRI, RBI);
}
return extractSubvector(ResVReg, ResType, ReadReg, Pos);
}
bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
Register ResourcePtr = I.getOperand(2).getReg();
SPIRVType *RegType = GR.getSPIRVTypeForVReg(ResourcePtr, I.getMF());
if (RegType->getOpcode() == SPIRV::OpTypeImage) {
// For texel buffers, the index into the image is part of the OpImageRead or
// OpImageWrite instructions. So we will do nothing in this case. This
// intrinsic will be combined with the load or store when selecting the load
// or store.
return true;
}
assert(ResType->getOpcode() == SPIRV::OpTypePointer);
MachineIRBuilder MIRBuilder(I);
Register IndexReg = I.getOperand(3).getReg();
Register ZeroReg =
buildZerosVal(GR.getOrCreateSPIRVIntegerType(32, I, TII), I);
return BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(ResourcePtr)
.addUse(ZeroReg)
.addUse(IndexReg)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectPushConstantGetPointer(
Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
MRI->replaceRegWith(ResVReg, I.getOperand(2).getReg());
return true;
}
bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
Register ObjReg = I.getOperand(2).getReg();
if (!BuildCOPY(ResVReg, ObjReg, I))
return false;
buildOpDecorate(ResVReg, I, TII, SPIRV::Decoration::NonUniformEXT, {});
// Check for the registers that use the index marked as non-uniform
// and recursively mark them as non-uniform.
// Per the spec, it's necessary that the final argument used for
// load/store/sample/atomic must be decorated, so we need to propagate the
// decoration through access chains and copies.
// https://docs.vulkan.org/samples/latest/samples/extensions/descriptor_indexing/README.html#_when_to_use_non_uniform_indexing_qualifier
decorateUsesAsNonUniform(ResVReg);
return true;
}
void SPIRVInstructionSelector::decorateUsesAsNonUniform(
Register &NonUniformReg) const {
llvm::SmallVector<Register> WorkList = {NonUniformReg};
while (WorkList.size() > 0) {
Register CurrentReg = WorkList.back();
WorkList.pop_back();
bool IsDecorated = false;
for (MachineInstr &Use : MRI->use_instructions(CurrentReg)) {
if (Use.getOpcode() == SPIRV::OpDecorate &&
Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
IsDecorated = true;
continue;
}
// Check if the instruction has the result register and add it to the
// worklist.
if (Use.getOperand(0).isReg() && Use.getOperand(0).isDef()) {
Register ResultReg = Use.getOperand(0).getReg();
if (ResultReg == CurrentReg)
continue;
WorkList.push_back(ResultReg);
}
}
if (!IsDecorated) {
buildOpDecorate(CurrentReg, *MRI->getVRegDef(CurrentReg), TII,
SPIRV::Decoration::NonUniformEXT, {});
}
}
}
bool SPIRVInstructionSelector::extractSubvector(
Register &ResVReg, const SPIRVType *ResType, Register &ReadReg,
MachineInstr &InsertionPoint) const {
SPIRVType *InputType = GR.getResultType(ReadReg);
[[maybe_unused]] uint64_t InputSize =
GR.getScalarOrVectorComponentCount(InputType);
uint64_t ResultSize = GR.getScalarOrVectorComponentCount(ResType);
assert(InputSize > 1 && "The input must be a vector.");
assert(ResultSize > 1 && "The result must be a vector.");
assert(ResultSize < InputSize &&
"Cannot extract more element than there are in the input.");
SmallVector<Register> ComponentRegisters;
SPIRVType *ScalarType = GR.getScalarOrVectorComponentType(ResType);
const TargetRegisterClass *ScalarRegClass = GR.getRegClass(ScalarType);
for (uint64_t I = 0; I < ResultSize; I++) {
Register ComponentReg = MRI->createVirtualRegister(ScalarRegClass);
bool Succeed = BuildMI(*InsertionPoint.getParent(), InsertionPoint,
InsertionPoint.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
.addDef(ComponentReg)
.addUse(ScalarType->getOperand(0).getReg())
.addUse(ReadReg)
.addImm(I)
.constrainAllUses(TII, TRI, RBI);
if (!Succeed)
return false;
ComponentRegisters.emplace_back(ComponentReg);
}
MachineInstrBuilder MIB = BuildMI(*InsertionPoint.getParent(), InsertionPoint,
InsertionPoint.getDebugLoc(),
TII.get(SPIRV::OpCompositeConstruct))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType));
for (Register ComponentReg : ComponentRegisters)
MIB.addUse(ComponentReg);
return MIB.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
MachineInstr &I) const {
// If the load of the image is in a different basic block, then
// this will generate invalid code. A proper solution is to move
// the OpLoad from selectHandleFromBinding here. However, to do
// that we will need to change the return type of the intrinsic.
// We will do that when we can, but for now trying to move forward with other
// issues.
Register ImageReg = I.getOperand(1).getReg();
auto *ImageDef = cast<GIntrinsic>(getVRegDef(*MRI, ImageReg));
Register NewImageReg = MRI->createVirtualRegister(MRI->getRegClass(ImageReg));
if (!loadHandleBeforePosition(NewImageReg, GR.getSPIRVTypeForVReg(ImageReg),
*ImageDef, I)) {
return false;
}
Register CoordinateReg = I.getOperand(2).getReg();
Register DataReg = I.getOperand(3).getReg();
assert(GR.getResultType(DataReg)->getOpcode() == SPIRV::OpTypeVector);
assert(GR.getScalarOrVectorComponentCount(GR.getResultType(DataReg)) == 4);
return BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
.addUse(NewImageReg)
.addUse(CoordinateReg)
.addUse(DataReg)
.constrainAllUses(TII, TRI, RBI);
}
Register SPIRVInstructionSelector::buildPointerToResource(
const SPIRVType *SpirvResType, SPIRV::StorageClass::StorageClass SC,
uint32_t Set, uint32_t Binding, uint32_t ArraySize, Register IndexReg,
StringRef Name, MachineIRBuilder MIRBuilder) const {
const Type *ResType = GR.getTypeForSPIRVType(SpirvResType);
if (ArraySize == 1) {
SPIRVType *PtrType =
GR.getOrCreateSPIRVPointerType(ResType, MIRBuilder, SC);
assert(GR.getPointeeType(PtrType) == SpirvResType &&
"SpirvResType did not have an explicit layout.");
return GR.getOrCreateGlobalVariableWithBinding(PtrType, Set, Binding, Name,
MIRBuilder);
}
const Type *VarType = ArrayType::get(const_cast<Type *>(ResType), ArraySize);
SPIRVType *VarPointerType =
GR.getOrCreateSPIRVPointerType(VarType, MIRBuilder, SC);
Register VarReg = GR.getOrCreateGlobalVariableWithBinding(
VarPointerType, Set, Binding, Name, MIRBuilder);
SPIRVType *ResPointerType =
GR.getOrCreateSPIRVPointerType(ResType, MIRBuilder, SC);
Register AcReg = MRI->createVirtualRegister(GR.getRegClass(ResPointerType));
MIRBuilder.buildInstr(SPIRV::OpAccessChain)
.addDef(AcReg)
.addUse(GR.getSPIRVTypeID(ResPointerType))
.addUse(VarReg)
.addUse(IndexReg);
return AcReg;
}
bool SPIRVInstructionSelector::selectFirstBitSet16(
Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
unsigned ExtendOpcode, unsigned BitSetOpcode) const {
Register ExtReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
bool Result = selectOpWithSrcs(ExtReg, ResType, I, {I.getOperand(2).getReg()},
ExtendOpcode);
return Result &&
selectFirstBitSet32(ResVReg, ResType, I, ExtReg, BitSetOpcode);
}
bool SPIRVInstructionSelector::selectFirstBitSet32(
Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
Register SrcReg, unsigned BitSetOpcode) const {
return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
.addImm(BitSetOpcode)
.addUse(SrcReg)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
Register SrcReg, unsigned BitSetOpcode, bool SwapPrimarySide) const {
// SPIR-V allow vectors of size 2,3,4 only. Calling with a larger vectors
// requires creating a param register and return register with an invalid
// vector size. If that is resolved, then this function can be used for
// vectors of any component size.
unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);
assert(ComponentCount < 5 && "Vec 5+ will generate invalid SPIR-V ops");
MachineIRBuilder MIRBuilder(I);
SPIRVType *BaseType = GR.retrieveScalarOrVectorIntType(ResType);
SPIRVType *I64Type = GR.getOrCreateSPIRVIntegerType(64, MIRBuilder);
SPIRVType *I64x2Type =
GR.getOrCreateSPIRVVectorType(I64Type, 2, MIRBuilder, false);
SPIRVType *Vec2ResType =
GR.getOrCreateSPIRVVectorType(BaseType, 2, MIRBuilder, false);
std::vector<Register> PartialRegs;
// Loops 0, 2, 4, ... but stops one loop early when ComponentCount is odd
unsigned CurrentComponent = 0;
for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
// This register holds the firstbitX result for each of the i64x2 vectors
// extracted from SrcReg
Register BitSetResult =
MRI->createVirtualRegister(GR.getRegClass(I64x2Type));
auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
.addDef(BitSetResult)
.addUse(GR.getSPIRVTypeID(I64x2Type))
.addUse(SrcReg)
.addUse(SrcReg)
.addImm(CurrentComponent)
.addImm(CurrentComponent + 1);
if (!MIB.constrainAllUses(TII, TRI, RBI))
return false;
Register SubVecBitSetReg =
MRI->createVirtualRegister(GR.getRegClass(Vec2ResType));
if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType, I, BitSetResult,
BitSetOpcode, SwapPrimarySide))
return false;
PartialRegs.push_back(SubVecBitSetReg);
}
// On odd component counts we need to handle one more component
if (CurrentComponent != ComponentCount) {
bool ZeroAsNull = !STI.isShader();
Register FinalElemReg = MRI->createVirtualRegister(GR.getRegClass(I64Type));
Register ConstIntLastIdx = GR.getOrCreateConstInt(
ComponentCount - 1, I, BaseType, TII, ZeroAsNull);
if (!selectOpWithSrcs(FinalElemReg, I64Type, I, {SrcReg, ConstIntLastIdx},
SPIRV::OpVectorExtractDynamic))
return false;
Register FinalElemBitSetReg =
MRI->createVirtualRegister(GR.getRegClass(BaseType));
if (!selectFirstBitSet64(FinalElemBitSetReg, BaseType, I, FinalElemReg,
BitSetOpcode, SwapPrimarySide))
return false;
PartialRegs.push_back(FinalElemBitSetReg);
}
// Join all the resulting registers back into the return type in order
// (ie i32x2, i32x2, i32x1 -> i32x5)
return selectOpWithSrcs(ResVReg, ResType, I, std::move(PartialRegs),
SPIRV::OpCompositeConstruct);
}
bool SPIRVInstructionSelector::selectFirstBitSet64(
Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
Register SrcReg, unsigned BitSetOpcode, bool SwapPrimarySide) const {
unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);
SPIRVType *BaseType = GR.retrieveScalarOrVectorIntType(ResType);
bool ZeroAsNull = !STI.isShader();
Register ConstIntZero =
GR.getOrCreateConstInt(0, I, BaseType, TII, ZeroAsNull);
Register ConstIntOne =
GR.getOrCreateConstInt(1, I, BaseType, TII, ZeroAsNull);
// SPIRV doesn't support vectors with more than 4 components. Since the
// algoritm below converts i64 -> i32x2 and i64x4 -> i32x8 it can only
// operate on vectors with 2 or less components. When largers vectors are
// seen. Split them, recurse, then recombine them.
if (ComponentCount > 2) {
return selectFirstBitSet64Overflow(ResVReg, ResType, I, SrcReg,
BitSetOpcode, SwapPrimarySide);
}
// 1. Split int64 into 2 pieces using a bitcast
MachineIRBuilder MIRBuilder(I);
SPIRVType *PostCastType = GR.getOrCreateSPIRVVectorType(
BaseType, 2 * ComponentCount, MIRBuilder, false);
Register BitcastReg =
MRI->createVirtualRegister(GR.getRegClass(PostCastType));
if (!selectOpWithSrcs(BitcastReg, PostCastType, I, {SrcReg},
SPIRV::OpBitcast))
return false;
// 2. Find the first set bit from the primary side for all the pieces in #1
Register FBSReg = MRI->createVirtualRegister(GR.getRegClass(PostCastType));
if (!selectFirstBitSet32(FBSReg, PostCastType, I, BitcastReg, BitSetOpcode))
return false;
// 3. Split result vector into high bits and low bits
Register HighReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
Register LowReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
bool IsScalarRes = ResType->getOpcode() != SPIRV::OpTypeVector;
if (IsScalarRes) {
// if scalar do a vector extract
if (!selectOpWithSrcs(HighReg, ResType, I, {FBSReg, ConstIntZero},
SPIRV::OpVectorExtractDynamic))
return false;
if (!selectOpWithSrcs(LowReg, ResType, I, {FBSReg, ConstIntOne},
SPIRV::OpVectorExtractDynamic))
return false;
} else {
// if vector do a shufflevector
auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
.addDef(HighReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(FBSReg)
// Per the spec, repeat the vector if only one vec is needed
.addUse(FBSReg);
// high bits are stored in even indexes. Extract them from FBSReg
for (unsigned J = 0; J < ComponentCount * 2; J += 2) {
MIB.addImm(J);
}
if (!MIB.constrainAllUses(TII, TRI, RBI))
return false;
MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
.addDef(LowReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(FBSReg)
// Per the spec, repeat the vector if only one vec is needed
.addUse(FBSReg);
// low bits are stored in odd indexes. Extract them from FBSReg
for (unsigned J = 1; J < ComponentCount * 2; J += 2) {
MIB.addImm(J);
}
if (!MIB.constrainAllUses(TII, TRI, RBI))
return false;
}
// 4. Check the result. When primary bits == -1 use secondary, otherwise use
// primary
SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
Register NegOneReg;
Register Reg0;
Register Reg32;
unsigned SelectOp;
unsigned AddOp;
if (IsScalarRes) {
NegOneReg =
GR.getOrCreateConstInt((unsigned)-1, I, ResType, TII, ZeroAsNull);
Reg0 = GR.getOrCreateConstInt(0, I, ResType, TII, ZeroAsNull);
Reg32 = GR.getOrCreateConstInt(32, I, ResType, TII, ZeroAsNull);
SelectOp = SPIRV::OpSelectSISCond;
AddOp = SPIRV::OpIAddS;
} else {
BoolType = GR.getOrCreateSPIRVVectorType(BoolType, ComponentCount,
MIRBuilder, false);
NegOneReg =
GR.getOrCreateConstVector((unsigned)-1, I, ResType, TII, ZeroAsNull);
Reg0 = GR.getOrCreateConstVector(0, I, ResType, TII, ZeroAsNull);
Reg32 = GR.getOrCreateConstVector(32, I, ResType, TII, ZeroAsNull);
SelectOp = SPIRV::OpSelectVIVCond;
AddOp = SPIRV::OpIAddV;
}
Register PrimaryReg = HighReg;
Register SecondaryReg = LowReg;
Register PrimaryShiftReg = Reg32;
Register SecondaryShiftReg = Reg0;
// By default the emitted opcodes check for the set bit from the MSB side.
// Setting SwapPrimarySide checks the set bit from the LSB side
if (SwapPrimarySide) {
PrimaryReg = LowReg;
SecondaryReg = HighReg;
PrimaryShiftReg = Reg0;
SecondaryShiftReg = Reg32;
}
// Check if the primary bits are == -1
Register BReg = MRI->createVirtualRegister(GR.getRegClass(BoolType));
if (!selectOpWithSrcs(BReg, BoolType, I, {PrimaryReg, NegOneReg},
SPIRV::OpIEqual))
return false;
// Select secondary bits if true in BReg, otherwise primary bits
Register TmpReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
if (!selectOpWithSrcs(TmpReg, ResType, I, {BReg, SecondaryReg, PrimaryReg},
SelectOp))
return false;
// 5. Add 32 when high bits are used, otherwise 0 for low bits
Register ValReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
if (!selectOpWithSrcs(ValReg, ResType, I,
{BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))
return false;
return selectOpWithSrcs(ResVReg, ResType, I, {ValReg, TmpReg}, AddOp);
}
bool SPIRVInstructionSelector::selectFirstBitHigh(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I,
bool IsSigned) const {
// FindUMsb and FindSMsb intrinsics only support 32 bit integers
Register OpReg = I.getOperand(2).getReg();
SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
// zero or sign extend
unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
switch (GR.getScalarOrVectorBitWidth(OpType)) {
case 16:
return selectFirstBitSet16(ResVReg, ResType, I, ExtendOpcode, BitSetOpcode);
case 32:
return selectFirstBitSet32(ResVReg, ResType, I, OpReg, BitSetOpcode);
case 64:
return selectFirstBitSet64(ResVReg, ResType, I, OpReg, BitSetOpcode,
/*SwapPrimarySide=*/false);
default:
report_fatal_error(
"spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
}
}
bool SPIRVInstructionSelector::selectFirstBitLow(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
// FindILsb intrinsic only supports 32 bit integers
Register OpReg = I.getOperand(2).getReg();
SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
// OpUConvert treats the operand bits as an unsigned i16 and zero extends it
// to an unsigned i32. As this leaves all the least significant bits unchanged
// so the first set bit from the LSB side doesn't change.
unsigned ExtendOpcode = SPIRV::OpUConvert;
unsigned BitSetOpcode = GL::FindILsb;
switch (GR.getScalarOrVectorBitWidth(OpType)) {
case 16:
return selectFirstBitSet16(ResVReg, ResType, I, ExtendOpcode, BitSetOpcode);
case 32:
return selectFirstBitSet32(ResVReg, ResType, I, OpReg, BitSetOpcode);
case 64:
return selectFirstBitSet64(ResVReg, ResType, I, OpReg, BitSetOpcode,
/*SwapPrimarySide=*/true);
default:
report_fatal_error("spv_firstbitlow only supports 16,32,64 bits.");
}
}
bool SPIRVInstructionSelector::selectAllocaArray(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
// there was an allocation size parameter to the allocation instruction
// that is not 1
MachineBasicBlock &BB = *I.getParent();
bool Res = BuildMI(BB, I, I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(I.getOperand(2).getReg())
.constrainAllUses(TII, TRI, RBI);
if (!STI.isShader()) {
unsigned Alignment = I.getOperand(3).getImm();
buildOpDecorate(ResVReg, I, TII, SPIRV::Decoration::Alignment, {Alignment});
}
return Res;
}
bool SPIRVInstructionSelector::selectFrameIndex(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
// Change order of instructions if needed: all OpVariable instructions in a
// function must be the first instructions in the first block
auto It = getOpVariableMBBIt(I);
bool Res = BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function))
.constrainAllUses(TII, TRI, RBI);
if (!STI.isShader()) {
unsigned Alignment = I.getOperand(2).getImm();
buildOpDecorate(ResVReg, *It, TII, SPIRV::Decoration::Alignment,
{Alignment});
}
return Res;
}
bool SPIRVInstructionSelector::selectBranch(MachineInstr &I) const {
// InstructionSelector walks backwards through the instructions. We can use
// both a G_BR and a G_BRCOND to create an OpBranchConditional. We hit G_BR
// first, so can generate an OpBranchConditional here. If there is no
// G_BRCOND, we just use OpBranch for a regular unconditional branch.
const MachineInstr *PrevI = I.getPrevNode();
MachineBasicBlock &MBB = *I.getParent();
if (PrevI != nullptr && PrevI->getOpcode() == TargetOpcode::G_BRCOND) {
return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranchConditional))
.addUse(PrevI->getOperand(0).getReg())
.addMBB(PrevI->getOperand(1).getMBB())
.addMBB(I.getOperand(0).getMBB())
.constrainAllUses(TII, TRI, RBI);
}
return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranch))
.addMBB(I.getOperand(0).getMBB())
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &I) const {
// InstructionSelector walks backwards through the instructions. For an
// explicit conditional branch with no fallthrough, we use both a G_BR and a
// G_BRCOND to create an OpBranchConditional. We should hit G_BR first, and
// generate the OpBranchConditional in selectBranch above.
//
// If an OpBranchConditional has been generated, we simply return, as the work
// is alread done. If there is no OpBranchConditional, LLVM must be relying on
// implicit fallthrough to the next basic block, so we need to create an
// OpBranchConditional with an explicit "false" argument pointing to the next
// basic block that LLVM would fall through to.
const MachineInstr *NextI = I.getNextNode();
// Check if this has already been successfully selected.
if (NextI != nullptr && NextI->getOpcode() == SPIRV::OpBranchConditional)
return true;
// Must be relying on implicit block fallthrough, so generate an
// OpBranchConditional with the "next" basic block as the "false" target.
MachineBasicBlock &MBB = *I.getParent();
unsigned NextMBBNum = MBB.getNextNode()->getNumber();
MachineBasicBlock *NextMBB = I.getMF()->getBlockNumbered(NextMBBNum);
return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranchConditional))
.addUse(I.getOperand(0).getReg())
.addMBB(I.getOperand(1).getMBB())
.addMBB(NextMBB)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectPhi(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpPhi))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType));
const unsigned NumOps = I.getNumOperands();
for (unsigned i = 1; i < NumOps; i += 2) {
MIB.addUse(I.getOperand(i + 0).getReg());
MIB.addMBB(I.getOperand(i + 1).getMBB());
}
bool Res = MIB.constrainAllUses(TII, TRI, RBI);
MIB->setDesc(TII.get(TargetOpcode::PHI));
MIB->removeOperand(1);
return Res;
}
bool SPIRVInstructionSelector::selectGlobalValue(
Register ResVReg, MachineInstr &I, const MachineInstr *Init) const {
// FIXME: don't use MachineIRBuilder here, replace it with BuildMI.
MachineIRBuilder MIRBuilder(I);
const GlobalValue *GV = I.getOperand(1).getGlobal();
Type *GVType = toTypedPointer(GR.getDeducedGlobalValueType(GV));
std::string GlobalIdent;
if (!GV->hasName()) {
unsigned &ID = UnnamedGlobalIDs[GV];
if (ID == 0)
ID = UnnamedGlobalIDs.size();
GlobalIdent = "__unnamed_" + Twine(ID).str();
} else {
GlobalIdent = GV->getName();
}
// Behaviour of functions as operands depends on availability of the
// corresponding extension (SPV_INTEL_function_pointers):
// - If there is an extension to operate with functions as operands:
// We create a proper constant operand and evaluate a correct type for a
// function pointer.
// - Without the required extension:
// We have functions as operands in tests with blocks of instruction e.g. in
// transcoding/global_block.ll. These operands are not used and should be
// substituted by zero constants. Their type is expected to be always
// OpTypePointer Function %uchar.
if (isa<Function>(GV)) {
const Constant *ConstVal = GV;
MachineBasicBlock &BB = *I.getParent();
Register NewReg = GR.find(ConstVal, GR.CurMF);
if (!NewReg.isValid()) {
Register NewReg = ResVReg;
const Function *GVFun =
STI.canUseExtension(SPIRV::Extension::SPV_INTEL_function_pointers)
? dyn_cast<Function>(GV)
: nullptr;
SPIRVType *ResType = GR.getOrCreateSPIRVPointerType(
GVType, I,
GVFun ? SPIRV::StorageClass::CodeSectionINTEL
: addressSpaceToStorageClass(GV->getAddressSpace(), STI));
if (GVFun) {
// References to a function via function pointers generate virtual
// registers without a definition. We will resolve it later, during
// module analysis stage.
Register ResTypeReg = GR.getSPIRVTypeID(ResType);
MachineRegisterInfo *MRI = MIRBuilder.getMRI();
Register FuncVReg =
MRI->createGenericVirtualRegister(GR.getRegType(ResType));
MRI->setRegClass(FuncVReg, &SPIRV::pIDRegClass);
MachineInstrBuilder MIB1 =
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
.addDef(FuncVReg)
.addUse(ResTypeReg);
MachineInstrBuilder MIB2 =
BuildMI(BB, I, I.getDebugLoc(),
TII.get(SPIRV::OpConstantFunctionPointerINTEL))
.addDef(NewReg)
.addUse(ResTypeReg)
.addUse(FuncVReg);
GR.add(ConstVal, MIB2);
// mapping the function pointer to the used Function
GR.recordFunctionPointer(&MIB2.getInstr()->getOperand(2), GVFun);
return MIB1.constrainAllUses(TII, TRI, RBI) &&
MIB2.constrainAllUses(TII, TRI, RBI);
}
MachineInstrBuilder MIB3 =
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
.addDef(NewReg)
.addUse(GR.getSPIRVTypeID(ResType));
GR.add(ConstVal, MIB3);
return MIB3.constrainAllUses(TII, TRI, RBI);
}
assert(NewReg != ResVReg);
return BuildCOPY(ResVReg, NewReg, I);
}
auto GlobalVar = cast<GlobalVariable>(GV);
assert(GlobalVar->getName() != "llvm.global.annotations");
// Skip empty declaration for GVs with initializers till we get the decl with
// passed initializer.
if (hasInitializer(GlobalVar) && !Init)
return true;
const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
getSpirvLinkageTypeFor(STI, *GV);
const unsigned AddrSpace = GV->getAddressSpace();
SPIRV::StorageClass::StorageClass StorageClass =
addressSpaceToStorageClass(AddrSpace, STI);
SPIRVType *ResType = GR.getOrCreateSPIRVPointerType(GVType, I, StorageClass);
Register Reg = GR.buildGlobalVariable(
ResVReg, ResType, GlobalIdent, GV, StorageClass, Init,
GlobalVar->isConstant(), LnkType, MIRBuilder, true);
return Reg.isValid();
}
bool SPIRVInstructionSelector::selectLog10(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
if (STI.canUseExtInstSet(SPIRV::InstructionSet::OpenCL_std)) {
return selectExtInst(ResVReg, ResType, I, CL::log10);
}
// There is no log10 instruction in the GLSL Extended Instruction set, so it
// is implemented as:
// log10(x) = log2(x) * (1 / log2(10))
// = log2(x) * 0.30103
MachineIRBuilder MIRBuilder(I);
MachineBasicBlock &BB = *I.getParent();
// Build log2(x).
Register VarReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
bool Result =
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
.addDef(VarReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
.addImm(GL::Log2)
.add(I.getOperand(1))
.constrainAllUses(TII, TRI, RBI);
// Build 0.30103.
assert(ResType->getOpcode() == SPIRV::OpTypeVector ||
ResType->getOpcode() == SPIRV::OpTypeFloat);
// TODO: Add matrix implementation once supported by the HLSL frontend.
const SPIRVType *SpirvScalarType =
ResType->getOpcode() == SPIRV::OpTypeVector
? GR.getSPIRVTypeForVReg(ResType->getOperand(1).getReg())
: ResType;
Register ScaleReg =
GR.buildConstantFP(APFloat(0.30103f), MIRBuilder, SpirvScalarType);
// Multiply log2(x) by 0.30103 to get log10(x) result.
auto Opcode = ResType->getOpcode() == SPIRV::OpTypeVector
? SPIRV::OpVectorTimesScalar
: SPIRV::OpFMulS;
return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(VarReg)
.addUse(ScaleReg)
.constrainAllUses(TII, TRI, RBI);
}
bool SPIRVInstructionSelector::selectModf(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
// llvm.modf has a single arg --the number to be decomposed-- and returns a
// struct { restype, restype }, while OpenCLLIB::modf has two args --the
// number to be decomposed and a pointer--, returns the fractional part and
// the integral part is stored in the pointer argument. Therefore, we can't
// use directly the OpenCLLIB::modf intrinsic. However, we can do some
// scaffolding to make it work. The idea is to create an alloca instruction
// to get a ptr, pass this ptr to OpenCL::modf, and then load the value
// from this ptr to place it in the struct. llvm.modf returns the fractional
// part as the first element of the result, and the integral part as the
// second element of the result.
// At this point, the return type is not a struct anymore, but rather two
// independent elements of SPIRVResType. We can get each independent element
// from I.getDefs() or I.getOperands().
if (STI.canUseExtInstSet(SPIRV::InstructionSet::OpenCL_std)) {
MachineIRBuilder MIRBuilder(I);
// Get pointer type for alloca variable.
const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(
ResType, MIRBuilder, SPIRV::StorageClass::Function);
// Create new register for the pointer type of alloca variable.
Register PtrTyReg =
MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::iIDRegClass);
MIRBuilder.getMRI()->setType(
PtrTyReg,
LLT::pointer(storageClassToAddressSpace(SPIRV::StorageClass::Function),
GR.getPointerSize()));
// Assign SPIR-V type of the pointer type of the alloca variable to the
// new register.
GR.assignSPIRVTypeToVReg(PtrType, PtrTyReg, MIRBuilder.getMF());
MachineBasicBlock &EntryBB = I.getMF()->front();
MachineBasicBlock::iterator VarPos =
getFirstValidInstructionInsertPoint(EntryBB);
auto AllocaMIB =
BuildMI(EntryBB, VarPos, I.getDebugLoc(), TII.get(SPIRV::OpVariable))
.addDef(PtrTyReg)
.addUse(GR.getSPIRVTypeID(PtrType))
.addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function));
Register Variable = AllocaMIB->getOperand(0).getReg();
MachineBasicBlock &BB = *I.getParent();
// Create the OpenCLLIB::modf instruction.
auto MIB =
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
.addImm(CL::modf)
.setMIFlags(I.getFlags())
.add(I.getOperand(I.getNumExplicitDefs())) // Floating point value.
.addUse(Variable); // Pointer to integral part.
// Assign the integral part stored in the ptr to the second element of the
// result.
Register IntegralPartReg = I.getOperand(1).getReg();
if (IntegralPartReg.isValid()) {
// Load the value from the pointer to integral part.
auto LoadMIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
.addDef(IntegralPartReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(Variable);
return LoadMIB.constrainAllUses(TII, TRI, RBI);
}
return MIB.constrainAllUses(TII, TRI, RBI);
} else if (STI.canUseExtInstSet(SPIRV::InstructionSet::GLSL_std_450)) {
assert(false && "GLSL::Modf is deprecated.");
// FIXME: GL::Modf is deprecated, use Modfstruct instead.
return false;
}
return false;
}
// Generate the instructions to load 3-element vector builtin input
// IDs/Indices.
// Like: GlobalInvocationId, LocalInvocationId, etc....
bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
SPIRV::BuiltIn::BuiltIn BuiltInValue, Register ResVReg,
const SPIRVType *ResType, MachineInstr &I) const {
MachineIRBuilder MIRBuilder(I);
const SPIRVType *Vec3Ty =
GR.getOrCreateSPIRVVectorType(ResType, 3, MIRBuilder, false);
const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(
Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
// Create new register for the input ID builtin variable.
Register NewRegister =
MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::iIDRegClass);
MIRBuilder.getMRI()->setType(NewRegister, LLT::pointer(0, 64));
GR.assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
// Build global variable with the necessary decorations for the input ID
// builtin variable.
Register Variable = GR.buildGlobalVariable(
NewRegister, PtrType, getLinkStringForBuiltIn(BuiltInValue), nullptr,
SPIRV::StorageClass::Input, nullptr, true, std::nullopt, MIRBuilder,
false);
// Create new register for loading value.
MachineRegisterInfo *MRI = MIRBuilder.getMRI();
Register LoadedRegister = MRI->createVirtualRegister(&SPIRV::iIDRegClass);
MIRBuilder.getMRI()->setType(LoadedRegister, LLT::pointer(0, 64));
GR.assignSPIRVTypeToVReg(Vec3Ty, LoadedRegister, MIRBuilder.getMF());
// Load v3uint value from the global variable.
bool Result =
BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
.addDef(LoadedRegister)
.addUse(GR.getSPIRVTypeID(Vec3Ty))
.addUse(Variable);
// Get the input ID index. Expecting operand is a constant immediate value,
// wrapped in a type assignment.
assert(I.getOperand(2).isReg());
const uint32_t ThreadId = foldImm(I.getOperand(2), MRI);
// Extract the input ID from the loaded vector value.
MachineBasicBlock &BB = *I.getParent();
auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(LoadedRegister)
.addImm(ThreadId);
return Result && MIB.constrainAllUses(TII, TRI, RBI);
}
// Generate the instructions to load 32-bit integer builtin input IDs/Indices.
// Like LocalInvocationIndex
bool SPIRVInstructionSelector::loadBuiltinInputID(
SPIRV::BuiltIn::BuiltIn BuiltInValue, Register ResVReg,
const SPIRVType *ResType, MachineInstr &I) const {
MachineIRBuilder MIRBuilder(I);
const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(
ResType, MIRBuilder, SPIRV::StorageClass::Input);
// Create new register for the input ID builtin variable.
Register NewRegister =
MIRBuilder.getMRI()->createVirtualRegister(GR.getRegClass(PtrType));
MIRBuilder.getMRI()->setType(
NewRegister,
LLT::pointer(storageClassToAddressSpace(SPIRV::StorageClass::Input),
GR.getPointerSize()));
GR.assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
// Build global variable with the necessary decorations for the input ID
// builtin variable.
Register Variable = GR.buildGlobalVariable(
NewRegister, PtrType, getLinkStringForBuiltIn(BuiltInValue), nullptr,
SPIRV::StorageClass::Input, nullptr, true, std::nullopt, MIRBuilder,
false);
// Load uint value from the global variable.
auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(Variable);
return MIB.constrainAllUses(TII, TRI, RBI);
}
SPIRVType *SPIRVInstructionSelector::widenTypeToVec4(const SPIRVType *Type,
MachineInstr &I) const {
MachineIRBuilder MIRBuilder(I);
if (Type->getOpcode() != SPIRV::OpTypeVector)
return GR.getOrCreateSPIRVVectorType(Type, 4, MIRBuilder, false);
uint64_t VectorSize = Type->getOperand(2).getImm();
if (VectorSize == 4)
return Type;
Register ScalarTypeReg = Type->getOperand(1).getReg();
const SPIRVType *ScalarType = GR.getSPIRVTypeForVReg(ScalarTypeReg);
return GR.getOrCreateSPIRVVectorType(ScalarType, 4, MIRBuilder, false);
}
bool SPIRVInstructionSelector::loadHandleBeforePosition(
Register &HandleReg, const SPIRVType *ResType, GIntrinsic &HandleDef,
MachineInstr &Pos) const {
assert(HandleDef.getIntrinsicID() ==
Intrinsic::spv_resource_handlefrombinding);
uint32_t Set = foldImm(HandleDef.getOperand(2), MRI);
uint32_t Binding = foldImm(HandleDef.getOperand(3), MRI);
uint32_t ArraySize = foldImm(HandleDef.getOperand(4), MRI);
Register IndexReg = HandleDef.getOperand(5).getReg();
std::string Name =
getStringValueFromReg(HandleDef.getOperand(6).getReg(), *MRI);
bool IsStructuredBuffer = ResType->getOpcode() == SPIRV::OpTypePointer;
MachineIRBuilder MIRBuilder(HandleDef);
SPIRVType *VarType = ResType;
SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
if (IsStructuredBuffer) {
VarType = GR.getPointeeType(ResType);
SC = GR.getPointerStorageClass(ResType);
}
Register VarReg = buildPointerToResource(VarType, SC, Set, Binding, ArraySize,
IndexReg, Name, MIRBuilder);
// The handle for the buffer is the pointer to the resource. For an image, the
// handle is the image object. So images get an extra load.
uint32_t LoadOpcode =
IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
GR.assignSPIRVTypeToVReg(ResType, HandleReg, *Pos.getMF());
return BuildMI(*Pos.getParent(), Pos, HandleDef.getDebugLoc(),
TII.get(LoadOpcode))
.addDef(HandleReg)
.addUse(GR.getSPIRVTypeID(ResType))
.addUse(VarReg)
.constrainAllUses(TII, TRI, RBI);
}
void SPIRVInstructionSelector::errorIfInstrOutsideShader(
MachineInstr &I) const {
if (!STI.isShader()) {
std::string DiagMsg;
raw_string_ostream OS(DiagMsg);
I.print(OS, true, false, false, false);
DiagMsg += " is only supported in shaders.\n";
report_fatal_error(DiagMsg.c_str(), false);
}
}
namespace llvm {
InstructionSelector *
createSPIRVInstructionSelector(const SPIRVTargetMachine &TM,
const SPIRVSubtarget &Subtarget,
const RegisterBankInfo &RBI) {
return new SPIRVInstructionSelector(TM, Subtarget, RBI);
}
} // namespace llvm
Reference in New Issue View Git Blame Copy Permalink