056f0a10b3
- [x] Implement refract using HLSL source in hlsl_intrinsics.h - [x] Implement the refract SPIR-V target built-in in clang/include/clang/Basic/BuiltinsSPIRV.td - [x] Add sema checks for refract to CheckSPIRVBuiltinFunctionCall in clang/lib/Sema/SemaSPIRV.cpp - [x] Add codegen for spv refract to EmitSPIRVBuiltinExpr in CGBuiltin.cpp - [x] Add codegen tests to clang/test/CodeGenHLSL/builtins/refract.hlsl - [x] Add spv codegen test to clang/test/CodeGenSPIRV/Builtins/refract.c - [x] Add sema tests to clang/test/SemaHLSL/BuiltIns/refract-errors.hlsl - [x] Add spv sema tests to clang/test/SemaSPIRV/BuiltIns/refract-errors.c - [x] Create the int_spv_refract intrinsic in IntrinsicsSPIRV.td - [x] In SPIRVInstructionSelector.cpp create the refract lowering and map it to int_spv_refract in SPIRVInstructionSelector::selectIntrinsic. - [x] Create SPIR-V backend test case in llvm/test/CodeGen/SPIRV/hlsl-intrinsics/refract.ll - [x] Check for what OpenCL support is needed. Resolves https://github.com/llvm/llvm-project/issues/99153
4177 lines
173 KiB
C++
4177 lines
173 KiB
C++
//===- 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);
|
|
|
|
// 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;
|
|
|
|
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 selectIntegerDot(Register ResVReg, const SPIRVType *ResType,
|
|
MachineInstr &I, bool Signed) const;
|
|
|
|
bool selectIntegerDotExpansion(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 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,
|
|
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 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 selectReadImageIntrinsic(Register &ResVReg, const SPIRVType *ResType,
|
|
MachineInstr &I) const;
|
|
bool selectImageWriteIntrinsic(MachineInstr &I) const;
|
|
bool selectResourceGetPointer(Register &ResVReg, const SPIRVType *ResType,
|
|
MachineInstr &I) 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,
|
|
bool IsNonUniform, 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 generateImageRead(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;
|
|
};
|
|
|
|
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()),
|
|
#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:
|
|
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;
|
|
}
|
|
|
|
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI) {
|
|
for (const auto &MO : MI.all_defs()) {
|
|
Register Reg = MO.getReg();
|
|
if (Reg.isPhysical() || !MRI.use_nodbg_empty(Reg))
|
|
return false;
|
|
}
|
|
if (MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE || MI.isFakeUse() ||
|
|
MI.isLifetimeMarker())
|
|
return false;
|
|
if (MI.isPHI())
|
|
return true;
|
|
if (MI.mayStore() || MI.isCall() ||
|
|
(MI.mayLoad() && MI.hasOrderedMemoryRef()) || MI.isPosition() ||
|
|
MI.isDebugInstr() || MI.isTerminator() || MI.isJumpTableDebugInfo())
|
|
return false;
|
|
return true;
|
|
}
|
|
|
|
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!");
|
|
|
|
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 = 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");
|
|
salvageDebugInfo(*MRI, I);
|
|
GR.invalidateMachineInstr(&I);
|
|
I.eraseFromParent();
|
|
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_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_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_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_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,
|
|
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_DEBUGTRAP:
|
|
case TargetOpcode::G_UBSANTRAP:
|
|
case TargetOpcode::DBG_LABEL:
|
|
return true;
|
|
|
|
default:
|
|
return false;
|
|
}
|
|
}
|
|
|
|
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);
|
|
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::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)) {
|
|
if ((*DefIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
|
|
(*DefIt).getOpcode() == SPIRV::OpVariable) {
|
|
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 generateImageRead(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);
|
|
}
|
|
|
|
bool SPIRVInstructionSelector::selectMemOperation(Register ResVReg,
|
|
MachineInstr &I) const {
|
|
MachineBasicBlock &BB = *I.getParent();
|
|
Register SrcReg = I.getOperand(1).getReg();
|
|
bool Result = true;
|
|
if (I.getOpcode() == TargetOpcode::G_MEMSET) {
|
|
MachineIRBuilder MIRBuilder(I);
|
|
assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());
|
|
unsigned Val = getIConstVal(I.getOperand(1).getReg(), MRI);
|
|
unsigned Num = getIConstVal(I.getOperand(2).getReg(), MRI);
|
|
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);
|
|
Register Const = GR.getOrCreateConstIntArray(Val, Num, I, SpvArrTy, TII);
|
|
// TODO: check if we have such GV, add init, use buildGlobalVariable.
|
|
Function &CurFunction = GR.CurMF->getFunction();
|
|
Type *LLVMArrTy =
|
|
ArrayType::get(IntegerType::get(CurFunction.getContext(), 8), Num);
|
|
// Module takes ownership of the global var.
|
|
GlobalVariable *GV = new GlobalVariable(*CurFunction.getParent(), LLVMArrTy,
|
|
true, GlobalValue::InternalLinkage,
|
|
Constant::getNullValue(LLVMArrTy));
|
|
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);
|
|
Result &= MIBVar.constrainAllUses(TII, TRI, RBI);
|
|
|
|
GR.add(GV, MIBVar);
|
|
GR.addGlobalObject(GV, GR.CurMF, VarReg);
|
|
|
|
buildOpDecorate(VarReg, I, TII, SPIRV::Decoration::Constant, {});
|
|
SPIRVType *SourceTy = GR.getOrCreateSPIRVPointerType(
|
|
ValTy, I, SPIRV::StorageClass::UniformConstant);
|
|
SrcReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
|
|
selectOpWithSrcs(SrcReg, SourceTy, I, {VarReg}, SPIRV::OpBitcast);
|
|
}
|
|
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);
|
|
}
|
|
Result &= MIB.constrainAllUses(TII, TRI, RBI);
|
|
if (ResVReg.isValid() && ResVReg != MIB->getOperand(0).getReg())
|
|
Result &= BuildCOPY(ResVReg, MIB->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 *DefType =
|
|
SrcReg.isValid() ? GR.getSPIRVTypeForVReg(SrcReg) : nullptr;
|
|
if (!DefType || DefType->getOpcode() != SPIRV::OpTypeVector)
|
|
report_fatal_error(
|
|
"cannot select G_UNMERGE_VALUES with a non-vector argument");
|
|
|
|
SPIRVType *ScalarType =
|
|
GR.getSPIRVTypeForVReg(DefType->getOperand(1).getReg());
|
|
MachineBasicBlock &BB = *I.getParent();
|
|
bool Res = false;
|
|
for (unsigned i = 0; i < I.getNumDefs(); ++i) {
|
|
Register ResVReg = I.getOperand(i).getReg();
|
|
SPIRVType *ResType = GR.getSPIRVTypeForVReg(ResVReg);
|
|
if (!ResType) {
|
|
// There was no "assign type" actions, let's fix this now
|
|
ResType = ScalarType;
|
|
MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
|
|
MRI->setType(ResVReg, LLT::scalar(GR.getScalarOrVectorBitWidth(ResType)));
|
|
GR.assignSPIRVTypeToVReg(ResType, ResVReg, *GR.CurMF);
|
|
}
|
|
auto MIB =
|
|
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
|
|
.addDef(ResVReg)
|
|
.addUse(GR.getSPIRVTypeID(ResType))
|
|
.addUse(SrcReg)
|
|
.addImm(static_cast<int64_t>(i));
|
|
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;
|
|
}
|
|
|
|
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::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)
|
|
.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,
|
|
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());
|
|
selectSelect(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 selectSelect(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());
|
|
// 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 ||
|
|
GR.isAggregateType(OpType)) {
|
|
// The case of G_ADDRSPACE_CAST inside spv_const_composite() is processed
|
|
// by selectAddrSpaceCast()
|
|
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::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_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_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_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_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_discard: {
|
|
return selectDiscard(ResVReg, ResType, I);
|
|
}
|
|
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::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 generateImageRead(ResVReg, ResType, NewImageReg, IdxReg, Loc, Pos);
|
|
}
|
|
|
|
bool SPIRVInstructionSelector::generateImageRead(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));
|
|
|
|
uint64_t ResultSize = GR.getScalarOrVectorComponentCount(ResType);
|
|
if (ResultSize == 4) {
|
|
auto BMI = BuildMI(*Pos.getParent(), Pos, Loc, TII.get(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(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::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,
|
|
bool IsNonUniform, 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));
|
|
if (IsNonUniform) {
|
|
// It is unclear which value needs to be marked an non-uniform, so both
|
|
// the index and the access changed are decorated as non-uniform.
|
|
buildOpDecorate(IndexReg, MIRBuilder, SPIRV::Decoration::NonUniformEXT, {});
|
|
buildOpDecorate(AcReg, MIRBuilder, SPIRV::Decoration::NonUniformEXT, {});
|
|
}
|
|
|
|
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, 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;
|
|
|
|
bool HasLnkTy = !GV->hasInternalLinkage() && !GV->hasPrivateLinkage() &&
|
|
!GV->hasHiddenVisibility();
|
|
SPIRV::LinkageType::LinkageType LnkType =
|
|
GV->isDeclarationForLinker()
|
|
? SPIRV::LinkageType::Import
|
|
: (GV->hasLinkOnceODRLinkage() &&
|
|
STI.canUseExtension(SPIRV::Extension::SPV_KHR_linkonce_odr)
|
|
? SPIRV::LinkageType::LinkOnceODR
|
|
: SPIRV::LinkageType::Export);
|
|
|
|
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(), HasLnkTy, 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);
|
|
}
|
|
|
|
// 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, false,
|
|
SPIRV::LinkageType::Import, 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, false,
|
|
SPIRV::LinkageType::Import, 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();
|
|
bool IsNonUniform = ArraySize > 1 && foldImm(HandleDef.getOperand(6), MRI);
|
|
std::string Name =
|
|
getStringValueFromReg(HandleDef.getOperand(7).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,
|
|
IsNonUniform, Name, MIRBuilder);
|
|
|
|
if (IsNonUniform)
|
|
buildOpDecorate(HandleReg, HandleDef, TII, SPIRV::Decoration::NonUniformEXT,
|
|
{});
|
|
|
|
// 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);
|
|
}
|
|
|
|
namespace llvm {
|
|
InstructionSelector *
|
|
createSPIRVInstructionSelector(const SPIRVTargetMachine &TM,
|
|
const SPIRVSubtarget &Subtarget,
|
|
const RegisterBankInfo &RBI) {
|
|
return new SPIRVInstructionSelector(TM, Subtarget, RBI);
|
|
}
|
|
} // namespace llvm
|