Google Git
Sign in
swiftshader / SwiftShader / 9c9608fa94a9209f2635c1d821c3652c3fd2a5e8 / . / third_party / llvm-16.0 / llvm / lib / Target / SPIRV / SPIRVBuiltins.cpp
blob: 2f44c999e5a22f69a5b4b4c0bf01f07a6cda5534 [file] [log] [blame]
//===- SPIRVBuiltins.cpp - SPIR-V Built-in Functions ------------*- 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 lowering builtin function calls and types using their
// demangled names and TableGen records.
//
//===----------------------------------------------------------------------===//
#include "SPIRVBuiltins.h"
#include "SPIRV.h"
#include "SPIRVUtils.h"
#include "llvm/Analysis/ValueTracking.h"
#include "llvm/IR/IntrinsicsSPIRV.h"
#include <string>
#include <tuple>
#define DEBUG_TYPE "spirv-builtins"
namespace llvm {
namespace SPIRV {
#define GET_BuiltinGroup_DECL
#include "SPIRVGenTables.inc"
struct DemangledBuiltin {
StringRef Name;
InstructionSet::InstructionSet Set;
BuiltinGroup Group;
uint8_t MinNumArgs;
uint8_t MaxNumArgs;
};
#define GET_DemangledBuiltins_DECL
#define GET_DemangledBuiltins_IMPL
struct IncomingCall {
const std::string BuiltinName;
const DemangledBuiltin *Builtin;
const Register ReturnRegister;
const SPIRVType *ReturnType;
const SmallVectorImpl<Register> &Arguments;
IncomingCall(const std::string BuiltinName, const DemangledBuiltin *Builtin,
const Register ReturnRegister, const SPIRVType *ReturnType,
const SmallVectorImpl<Register> &Arguments)
: BuiltinName(BuiltinName), Builtin(Builtin),
ReturnRegister(ReturnRegister), ReturnType(ReturnType),
Arguments(Arguments) {}
};
struct NativeBuiltin {
StringRef Name;
InstructionSet::InstructionSet Set;
uint32_t Opcode;
};
#define GET_NativeBuiltins_DECL
#define GET_NativeBuiltins_IMPL
struct GroupBuiltin {
StringRef Name;
uint32_t Opcode;
uint32_t GroupOperation;
bool IsElect;
bool IsAllOrAny;
bool IsAllEqual;
bool IsBallot;
bool IsInverseBallot;
bool IsBallotBitExtract;
bool IsBallotFindBit;
bool IsLogical;
bool NoGroupOperation;
bool HasBoolArg;
};
#define GET_GroupBuiltins_DECL
#define GET_GroupBuiltins_IMPL
struct GetBuiltin {
StringRef Name;
InstructionSet::InstructionSet Set;
BuiltIn::BuiltIn Value;
};
using namespace BuiltIn;
#define GET_GetBuiltins_DECL
#define GET_GetBuiltins_IMPL
struct ImageQueryBuiltin {
StringRef Name;
InstructionSet::InstructionSet Set;
uint32_t Component;
};
#define GET_ImageQueryBuiltins_DECL
#define GET_ImageQueryBuiltins_IMPL
struct ConvertBuiltin {
StringRef Name;
InstructionSet::InstructionSet Set;
bool IsDestinationSigned;
bool IsSaturated;
bool IsRounded;
FPRoundingMode::FPRoundingMode RoundingMode;
};
struct VectorLoadStoreBuiltin {
StringRef Name;
InstructionSet::InstructionSet Set;
uint32_t Number;
bool IsRounded;
FPRoundingMode::FPRoundingMode RoundingMode;
};
using namespace FPRoundingMode;
#define GET_ConvertBuiltins_DECL
#define GET_ConvertBuiltins_IMPL
using namespace InstructionSet;
#define GET_VectorLoadStoreBuiltins_DECL
#define GET_VectorLoadStoreBuiltins_IMPL
#define GET_CLMemoryScope_DECL
#define GET_CLSamplerAddressingMode_DECL
#define GET_CLMemoryFenceFlags_DECL
#define GET_ExtendedBuiltins_DECL
#include "SPIRVGenTables.inc"
} // namespace SPIRV
//===----------------------------------------------------------------------===//
// Misc functions for looking up builtins and veryfying requirements using
// TableGen records
//===----------------------------------------------------------------------===//
/// Looks up the demangled builtin call in the SPIRVBuiltins.td records using
/// the provided \p DemangledCall and specified \p Set.
///
/// The lookup follows the following algorithm, returning the first successful
/// match:
/// 1. Search with the plain demangled name (expecting a 1:1 match).
/// 2. Search with the prefix before or suffix after the demangled name
/// signyfying the type of the first argument.
///
/// \returns Wrapper around the demangled call and found builtin definition.
static std::unique_ptr<const SPIRV::IncomingCall>
lookupBuiltin(StringRef DemangledCall,
SPIRV::InstructionSet::InstructionSet Set,
Register ReturnRegister, const SPIRVType *ReturnType,
const SmallVectorImpl<Register> &Arguments) {
// Extract the builtin function name and types of arguments from the call
// skeleton.
std::string BuiltinName =
DemangledCall.substr(0, DemangledCall.find('(')).str();
// Check if the extracted name contains type information between angle
// brackets. If so, the builtin is an instantiated template - needs to have
// the information after angle brackets and return type removed.
if (BuiltinName.find('<') && BuiltinName.back() == '>') {
BuiltinName = BuiltinName.substr(0, BuiltinName.find('<'));
BuiltinName = BuiltinName.substr(BuiltinName.find_last_of(" ") + 1);
}
// Check if the extracted name begins with "__spirv_ImageSampleExplicitLod"
// contains return type information at the end "_R<type>", if so extract the
// plain builtin name without the type information.
if (StringRef(BuiltinName).contains("__spirv_ImageSampleExplicitLod") &&
StringRef(BuiltinName).contains("_R")) {
BuiltinName = BuiltinName.substr(0, BuiltinName.find("_R"));
}
SmallVector<StringRef, 10> BuiltinArgumentTypes;
StringRef BuiltinArgs =
DemangledCall.slice(DemangledCall.find('(') + 1, DemangledCall.find(')'));
BuiltinArgs.split(BuiltinArgumentTypes, ',', -1, false);
// Look up the builtin in the defined set. Start with the plain demangled
// name, expecting a 1:1 match in the defined builtin set.
const SPIRV::DemangledBuiltin *Builtin;
if ((Builtin = SPIRV::lookupBuiltin(BuiltinName, Set)))
return std::make_unique<SPIRV::IncomingCall>(
BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
// If the initial look up was unsuccessful and the demangled call takes at
// least 1 argument, add a prefix or suffix signifying the type of the first
// argument and repeat the search.
if (BuiltinArgumentTypes.size() >= 1) {
char FirstArgumentType = BuiltinArgumentTypes[0][0];
// Prefix to be added to the builtin's name for lookup.
// For example, OpenCL "abs" taking an unsigned value has a prefix "u_".
std::string Prefix;
switch (FirstArgumentType) {
// Unsigned:
case 'u':
if (Set == SPIRV::InstructionSet::OpenCL_std)
Prefix = "u_";
else if (Set == SPIRV::InstructionSet::GLSL_std_450)
Prefix = "u";
break;
// Signed:
case 'c':
case 's':
case 'i':
case 'l':
if (Set == SPIRV::InstructionSet::OpenCL_std)
Prefix = "s_";
else if (Set == SPIRV::InstructionSet::GLSL_std_450)
Prefix = "s";
break;
// Floating-point:
case 'f':
case 'd':
case 'h':
if (Set == SPIRV::InstructionSet::OpenCL_std ||
Set == SPIRV::InstructionSet::GLSL_std_450)
Prefix = "f";
break;
}
// If argument-type name prefix was added, look up the builtin again.
if (!Prefix.empty() &&
(Builtin = SPIRV::lookupBuiltin(Prefix + BuiltinName, Set)))
return std::make_unique<SPIRV::IncomingCall>(
BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
// If lookup with a prefix failed, find a suffix to be added to the
// builtin's name for lookup. For example, OpenCL "group_reduce_max" taking
// an unsigned value has a suffix "u".
std::string Suffix;
switch (FirstArgumentType) {
// Unsigned:
case 'u':
Suffix = "u";
break;
// Signed:
case 'c':
case 's':
case 'i':
case 'l':
Suffix = "s";
break;
// Floating-point:
case 'f':
case 'd':
case 'h':
Suffix = "f";
break;
}
// If argument-type name suffix was added, look up the builtin again.
if (!Suffix.empty() &&
(Builtin = SPIRV::lookupBuiltin(BuiltinName + Suffix, Set)))
return std::make_unique<SPIRV::IncomingCall>(
BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
}
// No builtin with such name was found in the set.
return nullptr;
}
//===----------------------------------------------------------------------===//
// Helper functions for building misc instructions
//===----------------------------------------------------------------------===//
/// Helper function building either a resulting scalar or vector bool register
/// depending on the expected \p ResultType.
///
/// \returns Tuple of the resulting register and its type.
static std::tuple<Register, SPIRVType *>
buildBoolRegister(MachineIRBuilder &MIRBuilder, const SPIRVType *ResultType,
SPIRVGlobalRegistry *GR) {
LLT Type;
SPIRVType *BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);
if (ResultType->getOpcode() == SPIRV::OpTypeVector) {
unsigned VectorElements = ResultType->getOperand(2).getImm();
BoolType =
GR->getOrCreateSPIRVVectorType(BoolType, VectorElements, MIRBuilder);
const FixedVectorType *LLVMVectorType =
cast<FixedVectorType>(GR->getTypeForSPIRVType(BoolType));
Type = LLT::vector(LLVMVectorType->getElementCount(), 1);
} else {
Type = LLT::scalar(1);
}
Register ResultRegister =
MIRBuilder.getMRI()->createGenericVirtualRegister(Type);
GR->assignSPIRVTypeToVReg(BoolType, ResultRegister, MIRBuilder.getMF());
return std::make_tuple(ResultRegister, BoolType);
}
/// Helper function for building either a vector or scalar select instruction
/// depending on the expected \p ResultType.
static bool buildSelectInst(MachineIRBuilder &MIRBuilder,
Register ReturnRegister, Register SourceRegister,
const SPIRVType *ReturnType,
SPIRVGlobalRegistry *GR) {
Register TrueConst, FalseConst;
if (ReturnType->getOpcode() == SPIRV::OpTypeVector) {
unsigned Bits = GR->getScalarOrVectorBitWidth(ReturnType);
uint64_t AllOnes = APInt::getAllOnesValue(Bits).getZExtValue();
TrueConst = GR->getOrCreateConsIntVector(AllOnes, MIRBuilder, ReturnType);
FalseConst = GR->getOrCreateConsIntVector(0, MIRBuilder, ReturnType);
} else {
TrueConst = GR->buildConstantInt(1, MIRBuilder, ReturnType);
FalseConst = GR->buildConstantInt(0, MIRBuilder, ReturnType);
}
return MIRBuilder.buildSelect(ReturnRegister, SourceRegister, TrueConst,
FalseConst);
}
/// Helper function for building a load instruction loading into the
/// \p DestinationReg.
static Register buildLoadInst(SPIRVType *BaseType, Register PtrRegister,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR, LLT LowLevelType,
Register DestinationReg = Register(0)) {
MachineRegisterInfo *MRI = MIRBuilder.getMRI();
if (!DestinationReg.isValid()) {
DestinationReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
MRI->setType(DestinationReg, LLT::scalar(32));
GR->assignSPIRVTypeToVReg(BaseType, DestinationReg, MIRBuilder.getMF());
}
// TODO: consider using correct address space and alignment (p0 is canonical
// type for selection though).
MachinePointerInfo PtrInfo = MachinePointerInfo();
MIRBuilder.buildLoad(DestinationReg, PtrRegister, PtrInfo, Align());
return DestinationReg;
}
/// Helper function for building a load instruction for loading a builtin global
/// variable of \p BuiltinValue value.
static Register buildBuiltinVariableLoad(MachineIRBuilder &MIRBuilder,
SPIRVType *VariableType,
SPIRVGlobalRegistry *GR,
SPIRV::BuiltIn::BuiltIn BuiltinValue,
LLT LLType,
Register Reg = Register(0)) {
Register NewRegister =
MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
MIRBuilder.getMRI()->setType(NewRegister,
LLT::pointer(0, GR->getPointerSize()));
SPIRVType *PtrType = GR->getOrCreateSPIRVPointerType(
VariableType, MIRBuilder, SPIRV::StorageClass::Input);
GR->assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
// Set up the global OpVariable with the necessary builtin decorations.
Register Variable = GR->buildGlobalVariable(
NewRegister, PtrType, getLinkStringForBuiltIn(BuiltinValue), nullptr,
SPIRV::StorageClass::Input, nullptr, true, true,
SPIRV::LinkageType::Import, MIRBuilder, false);
// Load the value from the global variable.
Register LoadedRegister =
buildLoadInst(VariableType, Variable, MIRBuilder, GR, LLType, Reg);
MIRBuilder.getMRI()->setType(LoadedRegister, LLType);
return LoadedRegister;
}
/// Helper external function for inserting ASSIGN_TYPE instuction between \p Reg
/// and its definition, set the new register as a destination of the definition,
/// assign SPIRVType to both registers. If SpirvTy is provided, use it as
/// SPIRVType in ASSIGN_TYPE, otherwise create it from \p Ty. Defined in
/// SPIRVPreLegalizer.cpp.
extern Register insertAssignInstr(Register Reg, Type *Ty, SPIRVType *SpirvTy,
SPIRVGlobalRegistry *GR,
MachineIRBuilder &MIB,
MachineRegisterInfo &MRI);
// TODO: Move to TableGen.
static SPIRV::MemorySemantics::MemorySemantics
getSPIRVMemSemantics(std::memory_order MemOrder) {
switch (MemOrder) {
case std::memory_order::memory_order_relaxed:
return SPIRV::MemorySemantics::None;
case std::memory_order::memory_order_acquire:
return SPIRV::MemorySemantics::Acquire;
case std::memory_order::memory_order_release:
return SPIRV::MemorySemantics::Release;
case std::memory_order::memory_order_acq_rel:
return SPIRV::MemorySemantics::AcquireRelease;
case std::memory_order::memory_order_seq_cst:
return SPIRV::MemorySemantics::SequentiallyConsistent;
default:
llvm_unreachable("Unknown CL memory scope");
}
}
static SPIRV::Scope::Scope getSPIRVScope(SPIRV::CLMemoryScope ClScope) {
switch (ClScope) {
case SPIRV::CLMemoryScope::memory_scope_work_item:
return SPIRV::Scope::Invocation;
case SPIRV::CLMemoryScope::memory_scope_work_group:
return SPIRV::Scope::Workgroup;
case SPIRV::CLMemoryScope::memory_scope_device:
return SPIRV::Scope::Device;
case SPIRV::CLMemoryScope::memory_scope_all_svm_devices:
return SPIRV::Scope::CrossDevice;
case SPIRV::CLMemoryScope::memory_scope_sub_group:
return SPIRV::Scope::Subgroup;
}
llvm_unreachable("Unknown CL memory scope");
}
static Register buildConstantIntReg(uint64_t Val, MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR,
unsigned BitWidth = 32) {
SPIRVType *IntType = GR->getOrCreateSPIRVIntegerType(BitWidth, MIRBuilder);
return GR->buildConstantInt(Val, MIRBuilder, IntType);
}
static Register buildScopeReg(Register CLScopeRegister,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR,
const MachineRegisterInfo *MRI) {
auto CLScope =
static_cast<SPIRV::CLMemoryScope>(getIConstVal(CLScopeRegister, MRI));
SPIRV::Scope::Scope Scope = getSPIRVScope(CLScope);
if (CLScope == static_cast<unsigned>(Scope))
return CLScopeRegister;
return buildConstantIntReg(Scope, MIRBuilder, GR);
}
static Register buildMemSemanticsReg(Register SemanticsRegister,
Register PtrRegister,
const MachineRegisterInfo *MRI,
SPIRVGlobalRegistry *GR) {
std::memory_order Order =
static_cast<std::memory_order>(getIConstVal(SemanticsRegister, MRI));
unsigned Semantics =
getSPIRVMemSemantics(Order) |
getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
if (Order == Semantics)
return SemanticsRegister;
return Register();
}
/// Helper function for translating atomic init to OpStore.
static bool buildAtomicInitInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder) {
assert(Call->Arguments.size() == 2 &&
"Need 2 arguments for atomic init translation");
MIRBuilder.buildInstr(SPIRV::OpStore)
.addUse(Call->Arguments[0])
.addUse(Call->Arguments[1]);
return true;
}
/// Helper function for building an atomic load instruction.
static bool buildAtomicLoadInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
Register PtrRegister = Call->Arguments[0];
// TODO: if true insert call to __translate_ocl_memory_sccope before
// OpAtomicLoad and the function implementation. We can use Translator's
// output for transcoding/atomic_explicit_arguments.cl as an example.
Register ScopeRegister;
if (Call->Arguments.size() > 1)
ScopeRegister = Call->Arguments[1];
else
ScopeRegister = buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
Register MemSemanticsReg;
if (Call->Arguments.size() > 2) {
// TODO: Insert call to __translate_ocl_memory_order before OpAtomicLoad.
MemSemanticsReg = Call->Arguments[2];
} else {
int Semantics =
SPIRV::MemorySemantics::SequentiallyConsistent |
getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
}
MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(PtrRegister)
.addUse(ScopeRegister)
.addUse(MemSemanticsReg);
return true;
}
/// Helper function for building an atomic store instruction.
static bool buildAtomicStoreInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
Register ScopeRegister =
buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
Register PtrRegister = Call->Arguments[0];
int Semantics =
SPIRV::MemorySemantics::SequentiallyConsistent |
getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
Register MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
.addUse(PtrRegister)
.addUse(ScopeRegister)
.addUse(MemSemanticsReg)
.addUse(Call->Arguments[1]);
return true;
}
/// Helper function for building an atomic compare-exchange instruction.
static bool buildAtomicCompareExchangeInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
unsigned Opcode =
SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
bool IsCmpxchg = Call->Builtin->Name.contains("cmpxchg");
MachineRegisterInfo *MRI = MIRBuilder.getMRI();
Register ObjectPtr = Call->Arguments[0]; // Pointer (volatile A *object.)
Register ExpectedArg = Call->Arguments[1]; // Comparator (C* expected).
Register Desired = Call->Arguments[2]; // Value (C Desired).
SPIRVType *SpvDesiredTy = GR->getSPIRVTypeForVReg(Desired);
LLT DesiredLLT = MRI->getType(Desired);
assert(GR->getSPIRVTypeForVReg(ObjectPtr)->getOpcode() ==
SPIRV::OpTypePointer);
unsigned ExpectedType = GR->getSPIRVTypeForVReg(ExpectedArg)->getOpcode();
assert(IsCmpxchg ? ExpectedType == SPIRV::OpTypeInt
: ExpectedType == SPIRV::OpTypePointer);
assert(GR->isScalarOfType(Desired, SPIRV::OpTypeInt));
SPIRVType *SpvObjectPtrTy = GR->getSPIRVTypeForVReg(ObjectPtr);
assert(SpvObjectPtrTy->getOperand(2).isReg() && "SPIRV type is expected");
auto StorageClass = static_cast<SPIRV::StorageClass::StorageClass>(
SpvObjectPtrTy->getOperand(1).getImm());
auto MemSemStorage = getMemSemanticsForStorageClass(StorageClass);
Register MemSemEqualReg;
Register MemSemUnequalReg;
uint64_t MemSemEqual =
IsCmpxchg
? SPIRV::MemorySemantics::None
: SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
uint64_t MemSemUnequal =
IsCmpxchg
? SPIRV::MemorySemantics::None
: SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
if (Call->Arguments.size() >= 4) {
assert(Call->Arguments.size() >= 5 &&
"Need 5+ args for explicit atomic cmpxchg");
auto MemOrdEq =
static_cast<std::memory_order>(getIConstVal(Call->Arguments[3], MRI));
auto MemOrdNeq =
static_cast<std::memory_order>(getIConstVal(Call->Arguments[4], MRI));
MemSemEqual = getSPIRVMemSemantics(MemOrdEq) | MemSemStorage;
MemSemUnequal = getSPIRVMemSemantics(MemOrdNeq) | MemSemStorage;
if (MemOrdEq == MemSemEqual)
MemSemEqualReg = Call->Arguments[3];
if (MemOrdNeq == MemSemEqual)
MemSemUnequalReg = Call->Arguments[4];
}
if (!MemSemEqualReg.isValid())
MemSemEqualReg = buildConstantIntReg(MemSemEqual, MIRBuilder, GR);
if (!MemSemUnequalReg.isValid())
MemSemUnequalReg = buildConstantIntReg(MemSemUnequal, MIRBuilder, GR);
Register ScopeReg;
auto Scope = IsCmpxchg ? SPIRV::Scope::Workgroup : SPIRV::Scope::Device;
if (Call->Arguments.size() >= 6) {
assert(Call->Arguments.size() == 6 &&
"Extra args for explicit atomic cmpxchg");
auto ClScope = static_cast<SPIRV::CLMemoryScope>(
getIConstVal(Call->Arguments[5], MRI));
Scope = getSPIRVScope(ClScope);
if (ClScope == static_cast<unsigned>(Scope))
ScopeReg = Call->Arguments[5];
}
if (!ScopeReg.isValid())
ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
Register Expected = IsCmpxchg
? ExpectedArg
: buildLoadInst(SpvDesiredTy, ExpectedArg, MIRBuilder,
GR, LLT::scalar(32));
MRI->setType(Expected, DesiredLLT);
Register Tmp = !IsCmpxchg ? MRI->createGenericVirtualRegister(DesiredLLT)
: Call->ReturnRegister;
GR->assignSPIRVTypeToVReg(SpvDesiredTy, Tmp, MIRBuilder.getMF());
SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
MIRBuilder.buildInstr(Opcode)
.addDef(Tmp)
.addUse(GR->getSPIRVTypeID(IntTy))
.addUse(ObjectPtr)
.addUse(ScopeReg)
.addUse(MemSemEqualReg)
.addUse(MemSemUnequalReg)
.addUse(Desired)
.addUse(Expected);
if (!IsCmpxchg) {
MIRBuilder.buildInstr(SPIRV::OpStore).addUse(ExpectedArg).addUse(Tmp);
MIRBuilder.buildICmp(CmpInst::ICMP_EQ, Call->ReturnRegister, Tmp, Expected);
}
return true;
}
/// Helper function for building an atomic load instruction.
static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;
Register ScopeRegister;
if (Call->Arguments.size() >= 4) {
assert(Call->Arguments.size() == 4 &&
"Too many args for explicit atomic RMW");
ScopeRegister = buildScopeReg(Call->Arguments[3], MIRBuilder, GR, MRI);
}
if (!ScopeRegister.isValid())
ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GR);
Register PtrRegister = Call->Arguments[0];
unsigned Semantics = SPIRV::MemorySemantics::None;
Register MemSemanticsReg;
if (Call->Arguments.size() >= 3)
MemSemanticsReg =
buildMemSemanticsReg(Call->Arguments[2], PtrRegister, MRI, GR);
if (!MemSemanticsReg.isValid())
MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
MIRBuilder.buildInstr(Opcode)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(PtrRegister)
.addUse(ScopeRegister)
.addUse(MemSemanticsReg)
.addUse(Call->Arguments[1]);
return true;
}
/// Helper function for building atomic flag instructions (e.g.
/// OpAtomicFlagTestAndSet).
static bool buildAtomicFlagInst(const SPIRV::IncomingCall *Call,
unsigned Opcode, MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
Register PtrRegister = Call->Arguments[0];
unsigned Semantics = SPIRV::MemorySemantics::SequentiallyConsistent;
Register MemSemanticsReg;
if (Call->Arguments.size() >= 2)
MemSemanticsReg =
buildMemSemanticsReg(Call->Arguments[1], PtrRegister, MRI, GR);
if (!MemSemanticsReg.isValid())
MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
assert((Opcode != SPIRV::OpAtomicFlagClear ||
(Semantics != SPIRV::MemorySemantics::Acquire &&
Semantics != SPIRV::MemorySemantics::AcquireRelease)) &&
"Invalid memory order argument!");
SPIRV::Scope::Scope Scope = SPIRV::Scope::Device;
Register ScopeRegister;
if (Call->Arguments.size() >= 3)
ScopeRegister = buildScopeReg(Call->Arguments[2], MIRBuilder, GR, MRI);
if (!ScopeRegister.isValid())
ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GR);
auto MIB = MIRBuilder.buildInstr(Opcode);
if (Opcode == SPIRV::OpAtomicFlagTestAndSet)
MIB.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType));
MIB.addUse(PtrRegister).addUse(ScopeRegister).addUse(MemSemanticsReg);
return true;
}
/// Helper function for building barriers, i.e., memory/control ordering
/// operations.
static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
unsigned MemFlags = getIConstVal(Call->Arguments[0], MRI);
unsigned MemSemantics = SPIRV::MemorySemantics::None;
if (MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE)
MemSemantics |= SPIRV::MemorySemantics::WorkgroupMemory;
if (MemFlags & SPIRV::CLK_GLOBAL_MEM_FENCE)
MemSemantics |= SPIRV::MemorySemantics::CrossWorkgroupMemory;
if (MemFlags & SPIRV::CLK_IMAGE_MEM_FENCE)
MemSemantics |= SPIRV::MemorySemantics::ImageMemory;
if (Opcode == SPIRV::OpMemoryBarrier) {
std::memory_order MemOrder =
static_cast<std::memory_order>(getIConstVal(Call->Arguments[1], MRI));
MemSemantics = getSPIRVMemSemantics(MemOrder) | MemSemantics;
} else {
MemSemantics |= SPIRV::MemorySemantics::SequentiallyConsistent;
}
Register MemSemanticsReg;
if (MemFlags == MemSemantics)
MemSemanticsReg = Call->Arguments[0];
else
MemSemanticsReg = buildConstantIntReg(MemSemantics, MIRBuilder, GR);
Register ScopeReg;
SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;
SPIRV::Scope::Scope MemScope = Scope;
if (Call->Arguments.size() >= 2) {
assert(
((Opcode != SPIRV::OpMemoryBarrier && Call->Arguments.size() == 2) ||
(Opcode == SPIRV::OpMemoryBarrier && Call->Arguments.size() == 3)) &&
"Extra args for explicitly scoped barrier");
Register ScopeArg = (Opcode == SPIRV::OpMemoryBarrier) ? Call->Arguments[2]
: Call->Arguments[1];
SPIRV::CLMemoryScope CLScope =
static_cast<SPIRV::CLMemoryScope>(getIConstVal(ScopeArg, MRI));
MemScope = getSPIRVScope(CLScope);
if (!(MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE) ||
(Opcode == SPIRV::OpMemoryBarrier))
Scope = MemScope;
if (CLScope == static_cast<unsigned>(Scope))
ScopeReg = Call->Arguments[1];
}
if (!ScopeReg.isValid())
ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
auto MIB = MIRBuilder.buildInstr(Opcode).addUse(ScopeReg);
if (Opcode != SPIRV::OpMemoryBarrier)
MIB.addUse(buildConstantIntReg(MemScope, MIRBuilder, GR));
MIB.addUse(MemSemanticsReg);
return true;
}
static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim) {
switch (dim) {
case SPIRV::Dim::DIM_1D:
case SPIRV::Dim::DIM_Buffer:
return 1;
case SPIRV::Dim::DIM_2D:
case SPIRV::Dim::DIM_Cube:
case SPIRV::Dim::DIM_Rect:
return 2;
case SPIRV::Dim::DIM_3D:
return 3;
default:
llvm_unreachable("Cannot get num components for given Dim");
}
}
/// Helper function for obtaining the number of size components.
static unsigned getNumSizeComponents(SPIRVType *imgType) {
assert(imgType->getOpcode() == SPIRV::OpTypeImage);
auto dim = static_cast<SPIRV::Dim::Dim>(imgType->getOperand(2).getImm());
unsigned numComps = getNumComponentsForDim(dim);
bool arrayed = imgType->getOperand(4).getImm() == 1;
return arrayed ? numComps + 1 : numComps;
}
//===----------------------------------------------------------------------===//
// Implementation functions for each builtin group
//===----------------------------------------------------------------------===//
static bool generateExtInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the extended instruction number in the TableGen records.
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
uint32_t Number =
SPIRV::lookupExtendedBuiltin(Builtin->Name, Builtin->Set)->Number;
// Build extended instruction.
auto MIB =
MIRBuilder.buildInstr(SPIRV::OpExtInst)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
.addImm(Number);
for (auto Argument : Call->Arguments)
MIB.addUse(Argument);
return true;
}
static bool generateRelationalInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the instruction opcode in the TableGen records.
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
unsigned Opcode =
SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
Register CompareRegister;
SPIRVType *RelationType;
std::tie(CompareRegister, RelationType) =
buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
// Build relational instruction.
auto MIB = MIRBuilder.buildInstr(Opcode)
.addDef(CompareRegister)
.addUse(GR->getSPIRVTypeID(RelationType));
for (auto Argument : Call->Arguments)
MIB.addUse(Argument);
// Build select instruction.
return buildSelectInst(MIRBuilder, Call->ReturnRegister, CompareRegister,
Call->ReturnType, GR);
}
static bool generateGroupInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
const SPIRV::GroupBuiltin *GroupBuiltin =
SPIRV::lookupGroupBuiltin(Builtin->Name);
const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
Register Arg0;
if (GroupBuiltin->HasBoolArg) {
Register ConstRegister = Call->Arguments[0];
auto ArgInstruction = getDefInstrMaybeConstant(ConstRegister, MRI);
// TODO: support non-constant bool values.
assert(ArgInstruction->getOpcode() == TargetOpcode::G_CONSTANT &&
"Only constant bool value args are supported");
if (GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode() !=
SPIRV::OpTypeBool)
Arg0 = GR->buildConstantInt(getIConstVal(ConstRegister, MRI), MIRBuilder,
GR->getOrCreateSPIRVBoolType(MIRBuilder));
}
Register GroupResultRegister = Call->ReturnRegister;
SPIRVType *GroupResultType = Call->ReturnType;
// TODO: maybe we need to check whether the result type is already boolean
// and in this case do not insert select instruction.
const bool HasBoolReturnTy =
GroupBuiltin->IsElect || GroupBuiltin->IsAllOrAny ||
GroupBuiltin->IsAllEqual || GroupBuiltin->IsLogical ||
GroupBuiltin->IsInverseBallot || GroupBuiltin->IsBallotBitExtract;
if (HasBoolReturnTy)
std::tie(GroupResultRegister, GroupResultType) =
buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
auto Scope = Builtin->Name.startswith("sub_group") ? SPIRV::Scope::Subgroup
: SPIRV::Scope::Workgroup;
Register ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GR);
// Build work/sub group instruction.
auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode)
.addDef(GroupResultRegister)
.addUse(GR->getSPIRVTypeID(GroupResultType))
.addUse(ScopeRegister);
if (!GroupBuiltin->NoGroupOperation)
MIB.addImm(GroupBuiltin->GroupOperation);
if (Call->Arguments.size() > 0) {
MIB.addUse(Arg0.isValid() ? Arg0 : Call->Arguments[0]);
for (unsigned i = 1; i < Call->Arguments.size(); i++)
MIB.addUse(Call->Arguments[i]);
}
// Build select instruction.
if (HasBoolReturnTy)
buildSelectInst(MIRBuilder, Call->ReturnRegister, GroupResultRegister,
Call->ReturnType, GR);
return true;
}
// These queries ask for a single size_t result for a given dimension index, e.g
// size_t get_global_id(uint dimindex). In SPIR-V, the builtins corresonding to
// these values are all vec3 types, so we need to extract the correct index or
// return defaultVal (0 or 1 depending on the query). We also handle extending
// or tuncating in case size_t does not match the expected result type's
// bitwidth.
//
// For a constant index >= 3 we generate:
// %res = OpConstant %SizeT 0
//
// For other indices we generate:
// %g = OpVariable %ptr_V3_SizeT Input
// OpDecorate %g BuiltIn XXX
// OpDecorate %g LinkageAttributes "__spirv_BuiltInXXX"
// OpDecorate %g Constant
// %loadedVec = OpLoad %V3_SizeT %g
//
// Then, if the index is constant < 3, we generate:
// %res = OpCompositeExtract %SizeT %loadedVec idx
// If the index is dynamic, we generate:
// %tmp = OpVectorExtractDynamic %SizeT %loadedVec %idx
// %cmp = OpULessThan %bool %idx %const_3
// %res = OpSelect %SizeT %cmp %tmp %const_0
//
// If the bitwidth of %res does not match the expected return type, we add an
// extend or truncate.
static bool genWorkgroupQuery(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR,
SPIRV::BuiltIn::BuiltIn BuiltinValue,
uint64_t DefaultValue) {
Register IndexRegister = Call->Arguments[0];
const unsigned ResultWidth = Call->ReturnType->getOperand(1).getImm();
const unsigned PointerSize = GR->getPointerSize();
const SPIRVType *PointerSizeType =
GR->getOrCreateSPIRVIntegerType(PointerSize, MIRBuilder);
MachineRegisterInfo *MRI = MIRBuilder.getMRI();
auto IndexInstruction = getDefInstrMaybeConstant(IndexRegister, MRI);
// Set up the final register to do truncation or extension on at the end.
Register ToTruncate = Call->ReturnRegister;
// If the index is constant, we can statically determine if it is in range.
bool IsConstantIndex =
IndexInstruction->getOpcode() == TargetOpcode::G_CONSTANT;
// If it's out of range (max dimension is 3), we can just return the constant
// default value (0 or 1 depending on which query function).
if (IsConstantIndex && getIConstVal(IndexRegister, MRI) >= 3) {
Register defaultReg = Call->ReturnRegister;
if (PointerSize != ResultWidth) {
defaultReg = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
GR->assignSPIRVTypeToVReg(PointerSizeType, defaultReg,
MIRBuilder.getMF());
ToTruncate = defaultReg;
}
auto NewRegister =
GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
MIRBuilder.buildCopy(defaultReg, NewRegister);
} else { // If it could be in range, we need to load from the given builtin.
auto Vec3Ty =
GR->getOrCreateSPIRVVectorType(PointerSizeType, 3, MIRBuilder);
Register LoadedVector =
buildBuiltinVariableLoad(MIRBuilder, Vec3Ty, GR, BuiltinValue,
LLT::fixed_vector(3, PointerSize));
// Set up the vreg to extract the result to (possibly a new temporary one).
Register Extracted = Call->ReturnRegister;
if (!IsConstantIndex || PointerSize != ResultWidth) {
Extracted = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
GR->assignSPIRVTypeToVReg(PointerSizeType, Extracted, MIRBuilder.getMF());
}
// Use Intrinsic::spv_extractelt so dynamic vs static extraction is
// handled later: extr = spv_extractelt LoadedVector, IndexRegister.
MachineInstrBuilder ExtractInst = MIRBuilder.buildIntrinsic(
Intrinsic::spv_extractelt, ArrayRef<Register>{Extracted}, true);
ExtractInst.addUse(LoadedVector).addUse(IndexRegister);
// If the index is dynamic, need check if it's < 3, and then use a select.
if (!IsConstantIndex) {
insertAssignInstr(Extracted, nullptr, PointerSizeType, GR, MIRBuilder,
*MRI);
auto IndexType = GR->getSPIRVTypeForVReg(IndexRegister);
auto BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);
Register CompareRegister =
MRI->createGenericVirtualRegister(LLT::scalar(1));
GR->assignSPIRVTypeToVReg(BoolType, CompareRegister, MIRBuilder.getMF());
// Use G_ICMP to check if idxVReg < 3.
MIRBuilder.buildICmp(CmpInst::ICMP_ULT, CompareRegister, IndexRegister,
GR->buildConstantInt(3, MIRBuilder, IndexType));
// Get constant for the default value (0 or 1 depending on which
// function).
Register DefaultRegister =
GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
// Get a register for the selection result (possibly a new temporary one).
Register SelectionResult = Call->ReturnRegister;
if (PointerSize != ResultWidth) {
SelectionResult =
MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
GR->assignSPIRVTypeToVReg(PointerSizeType, SelectionResult,
MIRBuilder.getMF());
}
// Create the final G_SELECT to return the extracted value or the default.
MIRBuilder.buildSelect(SelectionResult, CompareRegister, Extracted,
DefaultRegister);
ToTruncate = SelectionResult;
} else {
ToTruncate = Extracted;
}
}
// Alter the result's bitwidth if it does not match the SizeT value extracted.
if (PointerSize != ResultWidth)
MIRBuilder.buildZExtOrTrunc(Call->ReturnRegister, ToTruncate);
return true;
}
static bool generateBuiltinVar(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the builtin variable record.
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
SPIRV::BuiltIn::BuiltIn Value =
SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value;
if (Value == SPIRV::BuiltIn::GlobalInvocationId)
return genWorkgroupQuery(Call, MIRBuilder, GR, Value, 0);
// Build a load instruction for the builtin variable.
unsigned BitWidth = GR->getScalarOrVectorBitWidth(Call->ReturnType);
LLT LLType;
if (Call->ReturnType->getOpcode() == SPIRV::OpTypeVector)
LLType =
LLT::fixed_vector(Call->ReturnType->getOperand(2).getImm(), BitWidth);
else
LLType = LLT::scalar(BitWidth);
return buildBuiltinVariableLoad(MIRBuilder, Call->ReturnType, GR, Value,
LLType, Call->ReturnRegister);
}
static bool generateAtomicInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the instruction opcode in the TableGen records.
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
unsigned Opcode =
SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
switch (Opcode) {
case SPIRV::OpStore:
return buildAtomicInitInst(Call, MIRBuilder);
case SPIRV::OpAtomicLoad:
return buildAtomicLoadInst(Call, MIRBuilder, GR);
case SPIRV::OpAtomicStore:
return buildAtomicStoreInst(Call, MIRBuilder, GR);
case SPIRV::OpAtomicCompareExchange:
case SPIRV::OpAtomicCompareExchangeWeak:
return buildAtomicCompareExchangeInst(Call, MIRBuilder, GR);
case SPIRV::OpAtomicIAdd:
case SPIRV::OpAtomicISub:
case SPIRV::OpAtomicOr:
case SPIRV::OpAtomicXor:
case SPIRV::OpAtomicAnd:
case SPIRV::OpAtomicExchange:
return buildAtomicRMWInst(Call, Opcode, MIRBuilder, GR);
case SPIRV::OpMemoryBarrier:
return buildBarrierInst(Call, SPIRV::OpMemoryBarrier, MIRBuilder, GR);
case SPIRV::OpAtomicFlagTestAndSet:
case SPIRV::OpAtomicFlagClear:
return buildAtomicFlagInst(Call, Opcode, MIRBuilder, GR);
default:
return false;
}
}
static bool generateBarrierInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the instruction opcode in the TableGen records.
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
unsigned Opcode =
SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
return buildBarrierInst(Call, Opcode, MIRBuilder, GR);
}
static bool generateDotOrFMulInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
unsigned Opcode = GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode();
bool IsVec = Opcode == SPIRV::OpTypeVector;
// Use OpDot only in case of vector args and OpFMul in case of scalar args.
MIRBuilder.buildInstr(IsVec ? SPIRV::OpDot : SPIRV::OpFMulS)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(Call->Arguments[0])
.addUse(Call->Arguments[1]);
return true;
}
static bool generateGetQueryInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the builtin record.
SPIRV::BuiltIn::BuiltIn Value =
SPIRV::lookupGetBuiltin(Call->Builtin->Name, Call->Builtin->Set)->Value;
uint64_t IsDefault = (Value == SPIRV::BuiltIn::GlobalSize ||
Value == SPIRV::BuiltIn::WorkgroupSize ||
Value == SPIRV::BuiltIn::EnqueuedWorkgroupSize);
return genWorkgroupQuery(Call, MIRBuilder, GR, Value, IsDefault ? 1 : 0);
}
static bool generateImageSizeQueryInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the image size query component number in the TableGen records.
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
uint32_t Component =
SPIRV::lookupImageQueryBuiltin(Builtin->Name, Builtin->Set)->Component;
// Query result may either be a vector or a scalar. If return type is not a
// vector, expect only a single size component. Otherwise get the number of
// expected components.
SPIRVType *RetTy = Call->ReturnType;
unsigned NumExpectedRetComponents = RetTy->getOpcode() == SPIRV::OpTypeVector
? RetTy->getOperand(2).getImm()
: 1;
// Get the actual number of query result/size components.
SPIRVType *ImgType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
unsigned NumActualRetComponents = getNumSizeComponents(ImgType);
Register QueryResult = Call->ReturnRegister;
SPIRVType *QueryResultType = Call->ReturnType;
if (NumExpectedRetComponents != NumActualRetComponents) {
QueryResult = MIRBuilder.getMRI()->createGenericVirtualRegister(
LLT::fixed_vector(NumActualRetComponents, 32));
SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
QueryResultType = GR->getOrCreateSPIRVVectorType(
IntTy, NumActualRetComponents, MIRBuilder);
GR->assignSPIRVTypeToVReg(QueryResultType, QueryResult, MIRBuilder.getMF());
}
bool IsDimBuf = ImgType->getOperand(2).getImm() == SPIRV::Dim::DIM_Buffer;
unsigned Opcode =
IsDimBuf ? SPIRV::OpImageQuerySize : SPIRV::OpImageQuerySizeLod;
auto MIB = MIRBuilder.buildInstr(Opcode)
.addDef(QueryResult)
.addUse(GR->getSPIRVTypeID(QueryResultType))
.addUse(Call->Arguments[0]);
if (!IsDimBuf)
MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Lod id.
if (NumExpectedRetComponents == NumActualRetComponents)
return true;
if (NumExpectedRetComponents == 1) {
// Only 1 component is expected, build OpCompositeExtract instruction.
unsigned ExtractedComposite =
Component == 3 ? NumActualRetComponents - 1 : Component;
assert(ExtractedComposite < NumActualRetComponents &&
"Invalid composite index!");
MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(QueryResult)
.addImm(ExtractedComposite);
} else {
// More than 1 component is expected, fill a new vector.
auto MIB = MIRBuilder.buildInstr(SPIRV::OpVectorShuffle)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(QueryResult)
.addUse(QueryResult);
for (unsigned i = 0; i < NumExpectedRetComponents; ++i)
MIB.addImm(i < NumActualRetComponents ? i : 0xffffffff);
}
return true;
}
static bool generateImageMiscQueryInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt &&
"Image samples query result must be of int type!");
// Lookup the instruction opcode in the TableGen records.
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
unsigned Opcode =
SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
Register Image = Call->Arguments[0];
SPIRV::Dim::Dim ImageDimensionality = static_cast<SPIRV::Dim::Dim>(
GR->getSPIRVTypeForVReg(Image)->getOperand(2).getImm());
switch (Opcode) {
case SPIRV::OpImageQuerySamples:
assert(ImageDimensionality == SPIRV::Dim::DIM_2D &&
"Image must be of 2D dimensionality");
break;
case SPIRV::OpImageQueryLevels:
assert((ImageDimensionality == SPIRV::Dim::DIM_1D ||
ImageDimensionality == SPIRV::Dim::DIM_2D ||
ImageDimensionality == SPIRV::Dim::DIM_3D ||
ImageDimensionality == SPIRV::Dim::DIM_Cube) &&
"Image must be of 1D/2D/3D/Cube dimensionality");
break;
}
MIRBuilder.buildInstr(Opcode)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(Image);
return true;
}
// TODO: Move to TableGen.
static SPIRV::SamplerAddressingMode::SamplerAddressingMode
getSamplerAddressingModeFromBitmask(unsigned Bitmask) {
switch (Bitmask & SPIRV::CLK_ADDRESS_MODE_MASK) {
case SPIRV::CLK_ADDRESS_CLAMP:
return SPIRV::SamplerAddressingMode::Clamp;
case SPIRV::CLK_ADDRESS_CLAMP_TO_EDGE:
return SPIRV::SamplerAddressingMode::ClampToEdge;
case SPIRV::CLK_ADDRESS_REPEAT:
return SPIRV::SamplerAddressingMode::Repeat;
case SPIRV::CLK_ADDRESS_MIRRORED_REPEAT:
return SPIRV::SamplerAddressingMode::RepeatMirrored;
case SPIRV::CLK_ADDRESS_NONE:
return SPIRV::SamplerAddressingMode::None;
default:
llvm_unreachable("Unknown CL address mode");
}
}
static unsigned getSamplerParamFromBitmask(unsigned Bitmask) {
return (Bitmask & SPIRV::CLK_NORMALIZED_COORDS_TRUE) ? 1 : 0;
}
static SPIRV::SamplerFilterMode::SamplerFilterMode
getSamplerFilterModeFromBitmask(unsigned Bitmask) {
if (Bitmask & SPIRV::CLK_FILTER_LINEAR)
return SPIRV::SamplerFilterMode::Linear;
if (Bitmask & SPIRV::CLK_FILTER_NEAREST)
return SPIRV::SamplerFilterMode::Nearest;
return SPIRV::SamplerFilterMode::Nearest;
}
static bool generateReadImageInst(const StringRef DemangledCall,
const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
Register Image = Call->Arguments[0];
MachineRegisterInfo *MRI = MIRBuilder.getMRI();
if (DemangledCall.contains_insensitive("ocl_sampler")) {
Register Sampler = Call->Arguments[1];
if (!GR->isScalarOfType(Sampler, SPIRV::OpTypeSampler) &&
getDefInstrMaybeConstant(Sampler, MRI)->getOperand(1).isCImm()) {
uint64_t SamplerMask = getIConstVal(Sampler, MRI);
Sampler = GR->buildConstantSampler(
Register(), getSamplerAddressingModeFromBitmask(SamplerMask),
getSamplerParamFromBitmask(SamplerMask),
getSamplerFilterModeFromBitmask(SamplerMask), MIRBuilder,
GR->getSPIRVTypeForVReg(Sampler));
}
SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
SPIRVType *SampledImageType =
GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
Register SampledImage = MRI->createVirtualRegister(&SPIRV::IDRegClass);
MIRBuilder.buildInstr(SPIRV::OpSampledImage)
.addDef(SampledImage)
.addUse(GR->getSPIRVTypeID(SampledImageType))
.addUse(Image)
.addUse(Sampler);
Register Lod = GR->buildConstantFP(APFloat::getZero(APFloat::IEEEsingle()),
MIRBuilder);
SPIRVType *TempType = Call->ReturnType;
bool NeedsExtraction = false;
if (TempType->getOpcode() != SPIRV::OpTypeVector) {
TempType =
GR->getOrCreateSPIRVVectorType(Call->ReturnType, 4, MIRBuilder);
NeedsExtraction = true;
}
LLT LLType = LLT::scalar(GR->getScalarOrVectorBitWidth(TempType));
Register TempRegister = MRI->createGenericVirtualRegister(LLType);
GR->assignSPIRVTypeToVReg(TempType, TempRegister, MIRBuilder.getMF());
MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
.addDef(NeedsExtraction ? TempRegister : Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(TempType))
.addUse(SampledImage)
.addUse(Call->Arguments[2]) // Coordinate.
.addImm(SPIRV::ImageOperand::Lod)
.addUse(Lod);
if (NeedsExtraction)
MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(TempRegister)
.addImm(0);
} else if (DemangledCall.contains_insensitive("msaa")) {
MIRBuilder.buildInstr(SPIRV::OpImageRead)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(Image)
.addUse(Call->Arguments[1]) // Coordinate.
.addImm(SPIRV::ImageOperand::Sample)
.addUse(Call->Arguments[2]);
} else {
MIRBuilder.buildInstr(SPIRV::OpImageRead)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(Image)
.addUse(Call->Arguments[1]); // Coordinate.
}
return true;
}
static bool generateWriteImageInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
MIRBuilder.buildInstr(SPIRV::OpImageWrite)
.addUse(Call->Arguments[0]) // Image.
.addUse(Call->Arguments[1]) // Coordinate.
.addUse(Call->Arguments[2]); // Texel.
return true;
}
static bool generateSampleImageInst(const StringRef DemangledCall,
const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
if (Call->Builtin->Name.contains_insensitive(
"__translate_sampler_initializer")) {
// Build sampler literal.
uint64_t Bitmask = getIConstVal(Call->Arguments[0], MIRBuilder.getMRI());
Register Sampler = GR->buildConstantSampler(
Call->ReturnRegister, getSamplerAddressingModeFromBitmask(Bitmask),
getSamplerParamFromBitmask(Bitmask),
getSamplerFilterModeFromBitmask(Bitmask), MIRBuilder, Call->ReturnType);
return Sampler.isValid();
} else if (Call->Builtin->Name.contains_insensitive("__spirv_SampledImage")) {
// Create OpSampledImage.
Register Image = Call->Arguments[0];
SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
SPIRVType *SampledImageType =
GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
Register SampledImage =
Call->ReturnRegister.isValid()
? Call->ReturnRegister
: MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
MIRBuilder.buildInstr(SPIRV::OpSampledImage)
.addDef(SampledImage)
.addUse(GR->getSPIRVTypeID(SampledImageType))
.addUse(Image)
.addUse(Call->Arguments[1]); // Sampler.
return true;
} else if (Call->Builtin->Name.contains_insensitive(
"__spirv_ImageSampleExplicitLod")) {
// Sample an image using an explicit level of detail.
std::string ReturnType = DemangledCall.str();
if (DemangledCall.contains("_R")) {
ReturnType = ReturnType.substr(ReturnType.find("_R") + 2);
ReturnType = ReturnType.substr(0, ReturnType.find('('));
}
SPIRVType *Type = GR->getOrCreateSPIRVTypeByName(ReturnType, MIRBuilder);
MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Type))
.addUse(Call->Arguments[0]) // Image.
.addUse(Call->Arguments[1]) // Coordinate.
.addImm(SPIRV::ImageOperand::Lod)
.addUse(Call->Arguments[3]);
return true;
}
return false;
}
static bool generateSelectInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder) {
MIRBuilder.buildSelect(Call->ReturnRegister, Call->Arguments[0],
Call->Arguments[1], Call->Arguments[2]);
return true;
}
static bool generateSpecConstantInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the instruction opcode in the TableGen records.
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
unsigned Opcode =
SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
switch (Opcode) {
case SPIRV::OpSpecConstant: {
// Build the SpecID decoration.
unsigned SpecId =
static_cast<unsigned>(getIConstVal(Call->Arguments[0], MRI));
buildOpDecorate(Call->ReturnRegister, MIRBuilder, SPIRV::Decoration::SpecId,
{SpecId});
// Determine the constant MI.
Register ConstRegister = Call->Arguments[1];
const MachineInstr *Const = getDefInstrMaybeConstant(ConstRegister, MRI);
assert(Const &&
(Const->getOpcode() == TargetOpcode::G_CONSTANT ||
Const->getOpcode() == TargetOpcode::G_FCONSTANT) &&
"Argument should be either an int or floating-point constant");
// Determine the opcode and built the OpSpec MI.
const MachineOperand &ConstOperand = Const->getOperand(1);
if (Call->ReturnType->getOpcode() == SPIRV::OpTypeBool) {
assert(ConstOperand.isCImm() && "Int constant operand is expected");
Opcode = ConstOperand.getCImm()->getValue().getZExtValue()
? SPIRV::OpSpecConstantTrue
: SPIRV::OpSpecConstantFalse;
}
auto MIB = MIRBuilder.buildInstr(Opcode)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType));
if (Call->ReturnType->getOpcode() != SPIRV::OpTypeBool) {
if (Const->getOpcode() == TargetOpcode::G_CONSTANT)
addNumImm(ConstOperand.getCImm()->getValue(), MIB);
else
addNumImm(ConstOperand.getFPImm()->getValueAPF().bitcastToAPInt(), MIB);
}
return true;
}
case SPIRV::OpSpecConstantComposite: {
auto MIB = MIRBuilder.buildInstr(Opcode)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType));
for (unsigned i = 0; i < Call->Arguments.size(); i++)
MIB.addUse(Call->Arguments[i]);
return true;
}
default:
return false;
}
}
static MachineInstr *getBlockStructInstr(Register ParamReg,
MachineRegisterInfo *MRI) {
// We expect the following sequence of instructions:
// %0:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.alloca)
// or = G_GLOBAL_VALUE @block_literal_global
// %1:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.bitcast), %0
// %2:_(p4) = G_ADDRSPACE_CAST %1:_(pN)
MachineInstr *MI = MRI->getUniqueVRegDef(ParamReg);
assert(MI->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST &&
MI->getOperand(1).isReg());
Register BitcastReg = MI->getOperand(1).getReg();
MachineInstr *BitcastMI = MRI->getUniqueVRegDef(BitcastReg);
assert(isSpvIntrinsic(*BitcastMI, Intrinsic::spv_bitcast) &&
BitcastMI->getOperand(2).isReg());
Register ValueReg = BitcastMI->getOperand(2).getReg();
MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
return ValueMI;
}
// Return an integer constant corresponding to the given register and
// defined in spv_track_constant.
// TODO: maybe unify with prelegalizer pass.
static unsigned getConstFromIntrinsic(Register Reg, MachineRegisterInfo *MRI) {
MachineInstr *DefMI = MRI->getUniqueVRegDef(Reg);
assert(isSpvIntrinsic(*DefMI, Intrinsic::spv_track_constant) &&
DefMI->getOperand(2).isReg());
MachineInstr *DefMI2 = MRI->getUniqueVRegDef(DefMI->getOperand(2).getReg());
assert(DefMI2->getOpcode() == TargetOpcode::G_CONSTANT &&
DefMI2->getOperand(1).isCImm());
return DefMI2->getOperand(1).getCImm()->getValue().getZExtValue();
}
// Return type of the instruction result from spv_assign_type intrinsic.
// TODO: maybe unify with prelegalizer pass.
static const Type *getMachineInstrType(MachineInstr *MI) {
MachineInstr *NextMI = MI->getNextNode();
if (isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_name))
NextMI = NextMI->getNextNode();
Register ValueReg = MI->getOperand(0).getReg();
if (!isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_type) ||
NextMI->getOperand(1).getReg() != ValueReg)
return nullptr;
Type *Ty = getMDOperandAsType(NextMI->getOperand(2).getMetadata(), 0);
assert(Ty && "Type is expected");
return getTypedPtrEltType(Ty);
}
static const Type *getBlockStructType(Register ParamReg,
MachineRegisterInfo *MRI) {
// In principle, this information should be passed to us from Clang via
// an elementtype attribute. However, said attribute requires that
// the function call be an intrinsic, which is not. Instead, we rely on being
// able to trace this to the declaration of a variable: OpenCL C specification
// section 6.12.5 should guarantee that we can do this.
MachineInstr *MI = getBlockStructInstr(ParamReg, MRI);
if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
return getTypedPtrEltType(MI->getOperand(1).getGlobal()->getType());
assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&
"Blocks in OpenCL C must be traceable to allocation site");
return getMachineInstrType(MI);
}
// TODO: maybe move to the global register.
static SPIRVType *
getOrCreateSPIRVDeviceEventPointer(MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
Type *OpaqueType = StructType::getTypeByName(Context, "spirv.DeviceEvent");
if (!OpaqueType)
OpaqueType = StructType::getTypeByName(Context, "opencl.clk_event_t");
if (!OpaqueType)
OpaqueType = StructType::create(Context, "spirv.DeviceEvent");
unsigned SC0 = storageClassToAddressSpace(SPIRV::StorageClass::Function);
unsigned SC1 = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
Type *PtrType = PointerType::get(PointerType::get(OpaqueType, SC0), SC1);
return GR->getOrCreateSPIRVType(PtrType, MIRBuilder);
}
static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
MachineRegisterInfo *MRI = MIRBuilder.getMRI();
const DataLayout &DL = MIRBuilder.getDataLayout();
bool HasEvents = Call->Builtin->Name.find("events") != StringRef::npos;
const SPIRVType *Int32Ty = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
// Make vararg instructions before OpEnqueueKernel.
// Local sizes arguments: Sizes of block invoke arguments. Clang generates
// local size operands as an array, so we need to unpack them.
SmallVector<Register, 16> LocalSizes;
if (Call->Builtin->Name.find("_varargs") != StringRef::npos) {
const unsigned LocalSizeArrayIdx = HasEvents ? 9 : 6;
Register GepReg = Call->Arguments[LocalSizeArrayIdx];
MachineInstr *GepMI = MRI->getUniqueVRegDef(GepReg);
assert(isSpvIntrinsic(*GepMI, Intrinsic::spv_gep) &&
GepMI->getOperand(3).isReg());
Register ArrayReg = GepMI->getOperand(3).getReg();
MachineInstr *ArrayMI = MRI->getUniqueVRegDef(ArrayReg);
const Type *LocalSizeTy = getMachineInstrType(ArrayMI);
assert(LocalSizeTy && "Local size type is expected");
const uint64_t LocalSizeNum =
cast<ArrayType>(LocalSizeTy)->getNumElements();
unsigned SC = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
const LLT LLType = LLT::pointer(SC, GR->getPointerSize());
const SPIRVType *PointerSizeTy = GR->getOrCreateSPIRVPointerType(
Int32Ty, MIRBuilder, SPIRV::StorageClass::Function);
for (unsigned I = 0; I < LocalSizeNum; ++I) {
Register Reg =
MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
MIRBuilder.getMRI()->setType(Reg, LLType);
GR->assignSPIRVTypeToVReg(PointerSizeTy, Reg, MIRBuilder.getMF());
auto GEPInst = MIRBuilder.buildIntrinsic(Intrinsic::spv_gep,
ArrayRef<Register>{Reg}, true);
GEPInst
.addImm(GepMI->getOperand(2).getImm()) // In bound.
.addUse(ArrayMI->getOperand(0).getReg()) // Alloca.
.addUse(buildConstantIntReg(0, MIRBuilder, GR)) // Indices.
.addUse(buildConstantIntReg(I, MIRBuilder, GR));
LocalSizes.push_back(Reg);
}
}
// SPIRV OpEnqueueKernel instruction has 10+ arguments.
auto MIB = MIRBuilder.buildInstr(SPIRV::OpEnqueueKernel)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Int32Ty));
// Copy all arguments before block invoke function pointer.
const unsigned BlockFIdx = HasEvents ? 6 : 3;
for (unsigned i = 0; i < BlockFIdx; i++)
MIB.addUse(Call->Arguments[i]);
// If there are no event arguments in the original call, add dummy ones.
if (!HasEvents) {
MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Dummy num events.
Register NullPtr = GR->getOrCreateConstNullPtr(
MIRBuilder, getOrCreateSPIRVDeviceEventPointer(MIRBuilder, GR));
MIB.addUse(NullPtr); // Dummy wait events.
MIB.addUse(NullPtr); // Dummy ret event.
}
MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI);
assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
// Invoke: Pointer to invoke function.
MIB.addGlobalAddress(BlockMI->getOperand(1).getGlobal());
Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
// Param: Pointer to block literal.
MIB.addUse(BlockLiteralReg);
Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
// TODO: these numbers should be obtained from block literal structure.
// Param Size: Size of block literal structure.
MIB.addUse(buildConstantIntReg(DL.getTypeStoreSize(PType), MIRBuilder, GR));
// Param Aligment: Aligment of block literal structure.
MIB.addUse(
buildConstantIntReg(DL.getPrefTypeAlignment(PType), MIRBuilder, GR));
for (unsigned i = 0; i < LocalSizes.size(); i++)
MIB.addUse(LocalSizes[i]);
return true;
}
static bool generateEnqueueInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the instruction opcode in the TableGen records.
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
unsigned Opcode =
SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
switch (Opcode) {
case SPIRV::OpRetainEvent:
case SPIRV::OpReleaseEvent:
return MIRBuilder.buildInstr(Opcode).addUse(Call->Arguments[0]);
case SPIRV::OpCreateUserEvent:
case SPIRV::OpGetDefaultQueue:
return MIRBuilder.buildInstr(Opcode)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType));
case SPIRV::OpIsValidEvent:
return MIRBuilder.buildInstr(Opcode)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(Call->Arguments[0]);
case SPIRV::OpSetUserEventStatus:
return MIRBuilder.buildInstr(Opcode)
.addUse(Call->Arguments[0])
.addUse(Call->Arguments[1]);
case SPIRV::OpCaptureEventProfilingInfo:
return MIRBuilder.buildInstr(Opcode)
.addUse(Call->Arguments[0])
.addUse(Call->Arguments[1])
.addUse(Call->Arguments[2]);
case SPIRV::OpBuildNDRange: {
MachineRegisterInfo *MRI = MIRBuilder.getMRI();
SPIRVType *PtrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
assert(PtrType->getOpcode() == SPIRV::OpTypePointer &&
PtrType->getOperand(2).isReg());
Register TypeReg = PtrType->getOperand(2).getReg();
SPIRVType *StructType = GR->getSPIRVTypeForVReg(TypeReg);
Register TmpReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
GR->assignSPIRVTypeToVReg(StructType, TmpReg, MIRBuilder.getMF());
// Skip the first arg, it's the destination pointer. OpBuildNDRange takes
// three other arguments, so pass zero constant on absence.
unsigned NumArgs = Call->Arguments.size();
assert(NumArgs >= 2);
Register GlobalWorkSize = Call->Arguments[NumArgs < 4 ? 1 : 2];
Register LocalWorkSize =
NumArgs == 2 ? Register(0) : Call->Arguments[NumArgs < 4 ? 2 : 3];
Register GlobalWorkOffset = NumArgs <= 3 ? Register(0) : Call->Arguments[1];
if (NumArgs < 4) {
Register Const;
SPIRVType *SpvTy = GR->getSPIRVTypeForVReg(GlobalWorkSize);
if (SpvTy->getOpcode() == SPIRV::OpTypePointer) {
MachineInstr *DefInstr = MRI->getUniqueVRegDef(GlobalWorkSize);
assert(DefInstr && isSpvIntrinsic(*DefInstr, Intrinsic::spv_gep) &&
DefInstr->getOperand(3).isReg());
Register GWSPtr = DefInstr->getOperand(3).getReg();
// TODO: Maybe simplify generation of the type of the fields.
unsigned Size = Call->Builtin->Name.equals("ndrange_3D") ? 3 : 2;
unsigned BitWidth = GR->getPointerSize() == 64 ? 64 : 32;
Type *BaseTy = IntegerType::get(
MIRBuilder.getMF().getFunction().getContext(), BitWidth);
Type *FieldTy = ArrayType::get(BaseTy, Size);
SPIRVType *SpvFieldTy = GR->getOrCreateSPIRVType(FieldTy, MIRBuilder);
GlobalWorkSize = MRI->createVirtualRegister(&SPIRV::IDRegClass);
GR->assignSPIRVTypeToVReg(SpvFieldTy, GlobalWorkSize,
MIRBuilder.getMF());
MIRBuilder.buildInstr(SPIRV::OpLoad)
.addDef(GlobalWorkSize)
.addUse(GR->getSPIRVTypeID(SpvFieldTy))
.addUse(GWSPtr);
Const = GR->getOrCreateConsIntArray(0, MIRBuilder, SpvFieldTy);
} else {
Const = GR->buildConstantInt(0, MIRBuilder, SpvTy);
}
if (!LocalWorkSize.isValid())
LocalWorkSize = Const;
if (!GlobalWorkOffset.isValid())
GlobalWorkOffset = Const;
}
MIRBuilder.buildInstr(Opcode)
.addDef(TmpReg)
.addUse(TypeReg)
.addUse(GlobalWorkSize)
.addUse(LocalWorkSize)
.addUse(GlobalWorkOffset);
return MIRBuilder.buildInstr(SPIRV::OpStore)
.addUse(Call->Arguments[0])
.addUse(TmpReg);
}
case SPIRV::OpEnqueueKernel:
return buildEnqueueKernel(Call, MIRBuilder, GR);
default:
return false;
}
}
static bool generateAsyncCopy(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the instruction opcode in the TableGen records.
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
unsigned Opcode =
SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
auto Scope = buildConstantIntReg(SPIRV::Scope::Workgroup, MIRBuilder, GR);
switch (Opcode) {
case SPIRV::OpGroupAsyncCopy:
return MIRBuilder.buildInstr(Opcode)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(Scope)
.addUse(Call->Arguments[0])
.addUse(Call->Arguments[1])
.addUse(Call->Arguments[2])
.addUse(buildConstantIntReg(1, MIRBuilder, GR))
.addUse(Call->Arguments[3]);
case SPIRV::OpGroupWaitEvents:
return MIRBuilder.buildInstr(Opcode)
.addUse(Scope)
.addUse(Call->Arguments[0])
.addUse(Call->Arguments[1]);
default:
return false;
}
}
static bool generateConvertInst(const StringRef DemangledCall,
const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the conversion builtin in the TableGen records.
const SPIRV::ConvertBuiltin *Builtin =
SPIRV::lookupConvertBuiltin(Call->Builtin->Name, Call->Builtin->Set);
if (Builtin->IsSaturated)
buildOpDecorate(Call->ReturnRegister, MIRBuilder,
SPIRV::Decoration::SaturatedConversion, {});
if (Builtin->IsRounded)
buildOpDecorate(Call->ReturnRegister, MIRBuilder,
SPIRV::Decoration::FPRoundingMode, {Builtin->RoundingMode});
unsigned Opcode = SPIRV::OpNop;
if (GR->isScalarOrVectorOfType(Call->Arguments[0], SPIRV::OpTypeInt)) {
// Int -> ...
if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {
// Int -> Int
if (Builtin->IsSaturated)
Opcode = Builtin->IsDestinationSigned ? SPIRV::OpSatConvertUToS
: SPIRV::OpSatConvertSToU;
else
Opcode = Builtin->IsDestinationSigned ? SPIRV::OpUConvert
: SPIRV::OpSConvert;
} else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
SPIRV::OpTypeFloat)) {
// Int -> Float
bool IsSourceSigned =
DemangledCall[DemangledCall.find_first_of('(') + 1] != 'u';
Opcode = IsSourceSigned ? SPIRV::OpConvertSToF : SPIRV::OpConvertUToF;
}
} else if (GR->isScalarOrVectorOfType(Call->Arguments[0],
SPIRV::OpTypeFloat)) {
// Float -> ...
if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt))
// Float -> Int
Opcode = Builtin->IsDestinationSigned ? SPIRV::OpConvertFToS
: SPIRV::OpConvertFToU;
else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
SPIRV::OpTypeFloat))
// Float -> Float
Opcode = SPIRV::OpFConvert;
}
assert(Opcode != SPIRV::OpNop &&
"Conversion between the types not implemented!");
MIRBuilder.buildInstr(Opcode)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addUse(Call->Arguments[0]);
return true;
}
static bool generateVectorLoadStoreInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the vector load/store builtin in the TableGen records.
const SPIRV::VectorLoadStoreBuiltin *Builtin =
SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name,
Call->Builtin->Set);
// Build extended instruction.
auto MIB =
MIRBuilder.buildInstr(SPIRV::OpExtInst)
.addDef(Call->ReturnRegister)
.addUse(GR->getSPIRVTypeID(Call->ReturnType))
.addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
.addImm(Builtin->Number);
for (auto Argument : Call->Arguments)
MIB.addUse(Argument);
// Rounding mode should be passed as a last argument in the MI for builtins
// like "vstorea_halfn_r".
if (Builtin->IsRounded)
MIB.addImm(static_cast<uint32_t>(Builtin->RoundingMode));
return true;
}
static bool generateLoadStoreInst(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup the instruction opcode in the TableGen records.
const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
unsigned Opcode =
SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
bool IsLoad = Opcode == SPIRV::OpLoad;
// Build the instruction.
auto MIB = MIRBuilder.buildInstr(Opcode);
if (IsLoad) {
MIB.addDef(Call->ReturnRegister);
MIB.addUse(GR->getSPIRVTypeID(Call->ReturnType));
}
// Add a pointer to the value to load/store.
MIB.addUse(Call->Arguments[0]);
// Add a value to store.
if (!IsLoad)
MIB.addUse(Call->Arguments[1]);
// Add optional memory attributes and an alignment.
MachineRegisterInfo *MRI = MIRBuilder.getMRI();
unsigned NumArgs = Call->Arguments.size();
if ((IsLoad && NumArgs >= 2) || NumArgs >= 3)
MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 1 : 2], MRI));
if ((IsLoad && NumArgs >= 3) || NumArgs >= 4)
MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 2 : 3], MRI));
return true;
}
/// Lowers a builtin funtion call using the provided \p DemangledCall skeleton
/// and external instruction \p Set.
namespace SPIRV {
std::optional<bool> lowerBuiltin(const StringRef DemangledCall,
SPIRV::InstructionSet::InstructionSet Set,
MachineIRBuilder &MIRBuilder,
const Register OrigRet, const Type *OrigRetTy,
const SmallVectorImpl<Register> &Args,
SPIRVGlobalRegistry *GR) {
LLVM_DEBUG(dbgs() << "Lowering builtin call: " << DemangledCall << "\n");
// SPIR-V type and return register.
Register ReturnRegister = OrigRet;
SPIRVType *ReturnType = nullptr;
if (OrigRetTy && !OrigRetTy->isVoidTy()) {
ReturnType = GR->assignTypeToVReg(OrigRetTy, OrigRet, MIRBuilder);
} else if (OrigRetTy && OrigRetTy->isVoidTy()) {
ReturnRegister = MIRBuilder.getMRI()->createVirtualRegister(&IDRegClass);
MIRBuilder.getMRI()->setType(ReturnRegister, LLT::scalar(32));
ReturnType = GR->assignTypeToVReg(OrigRetTy, ReturnRegister, MIRBuilder);
}
// Lookup the builtin in the TableGen records.
std::unique_ptr<const IncomingCall> Call =
lookupBuiltin(DemangledCall, Set, ReturnRegister, ReturnType, Args);
if (!Call) {
LLVM_DEBUG(dbgs() << "Builtin record was not found!\n");
return std::nullopt;
}
// TODO: check if the provided args meet the builtin requirments.
assert(Args.size() >= Call->Builtin->MinNumArgs &&
"Too few arguments to generate the builtin");
if (Call->Builtin->MaxNumArgs && Args.size() > Call->Builtin->MaxNumArgs)
LLVM_DEBUG(dbgs() << "More arguments provided than required!\n");
// Match the builtin with implementation based on the grouping.
switch (Call->Builtin->Group) {
case SPIRV::Extended:
return generateExtInst(Call.get(), MIRBuilder, GR);
case SPIRV::Relational:
return generateRelationalInst(Call.get(), MIRBuilder, GR);
case SPIRV::Group:
return generateGroupInst(Call.get(), MIRBuilder, GR);
case SPIRV::Variable:
return generateBuiltinVar(Call.get(), MIRBuilder, GR);
case SPIRV::Atomic:
return generateAtomicInst(Call.get(), MIRBuilder, GR);
case SPIRV::Barrier:
return generateBarrierInst(Call.get(), MIRBuilder, GR);
case SPIRV::Dot:
return generateDotOrFMulInst(Call.get(), MIRBuilder, GR);
case SPIRV::GetQuery:
return generateGetQueryInst(Call.get(), MIRBuilder, GR);
case SPIRV::ImageSizeQuery:
return generateImageSizeQueryInst(Call.get(), MIRBuilder, GR);
case SPIRV::ImageMiscQuery:
return generateImageMiscQueryInst(Call.get(), MIRBuilder, GR);
case SPIRV::ReadImage:
return generateReadImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
case SPIRV::WriteImage:
return generateWriteImageInst(Call.get(), MIRBuilder, GR);
case SPIRV::SampleImage:
return generateSampleImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
case SPIRV::Select:
return generateSelectInst(Call.get(), MIRBuilder);
case SPIRV::SpecConstant:
return generateSpecConstantInst(Call.get(), MIRBuilder, GR);
case SPIRV::Enqueue:
return generateEnqueueInst(Call.get(), MIRBuilder, GR);
case SPIRV::AsyncCopy:
return generateAsyncCopy(Call.get(), MIRBuilder, GR);
case SPIRV::Convert:
return generateConvertInst(DemangledCall, Call.get(), MIRBuilder, GR);
case SPIRV::VectorLoadStore:
return generateVectorLoadStoreInst(Call.get(), MIRBuilder, GR);
case SPIRV::LoadStore:
return generateLoadStoreInst(Call.get(), MIRBuilder, GR);
}
return false;
}
struct DemangledType {
StringRef Name;
uint32_t Opcode;
};
#define GET_DemangledTypes_DECL
#define GET_DemangledTypes_IMPL
struct ImageType {
StringRef Name;
StringRef SampledType;
AccessQualifier::AccessQualifier Qualifier;
Dim::Dim Dimensionality;
bool Arrayed;
bool Depth;
bool Multisampled;
bool Sampled;
ImageFormat::ImageFormat Format;
};
struct PipeType {
StringRef Name;
AccessQualifier::AccessQualifier Qualifier;
};
using namespace AccessQualifier;
using namespace Dim;
using namespace ImageFormat;
#define GET_ImageTypes_DECL
#define GET_ImageTypes_IMPL
#define GET_PipeTypes_DECL
#define GET_PipeTypes_IMPL
#include "SPIRVGenTables.inc"
} // namespace SPIRV
//===----------------------------------------------------------------------===//
// Misc functions for parsing builtin types and looking up implementation
// details in TableGenerated tables.
//===----------------------------------------------------------------------===//
static const SPIRV::DemangledType *findBuiltinType(StringRef Name) {
if (Name.startswith("opencl."))
return SPIRV::lookupBuiltinType(Name);
if (!Name.startswith("spirv."))
return nullptr;
// Some SPIR-V builtin types have a complex list of parameters as part of
// their name (e.g. spirv.Image._void_1_0_0_0_0_0_0). Those parameters often
// are numeric literals which cannot be easily represented by TableGen
// records and should be parsed instead.
unsigned BaseTypeNameLength =
Name.contains('_') ? Name.find('_') - 1 : Name.size();
return SPIRV::lookupBuiltinType(Name.substr(0, BaseTypeNameLength).str());
}
static std::unique_ptr<const SPIRV::ImageType>
lookupOrParseBuiltinImageType(StringRef Name) {
if (Name.startswith("opencl.")) {
// Lookup OpenCL builtin image type lowering details in TableGen records.
const SPIRV::ImageType *Record = SPIRV::lookupImageType(Name);
return std::unique_ptr<SPIRV::ImageType>(new SPIRV::ImageType(*Record));
}
if (!Name.startswith("spirv."))
llvm_unreachable("Unknown builtin image type name/literal");
// Parse the literals of SPIR-V image builtin parameters. The name should
// have the following format:
// spirv.Image._Type_Dim_Depth_Arrayed_MS_Sampled_ImageFormat_AccessQualifier
// e.g. %spirv.Image._void_1_0_0_0_0_0_0
StringRef TypeParametersString = Name.substr(strlen("spirv.Image."));
SmallVector<StringRef> TypeParameters;
SplitString(TypeParametersString, TypeParameters, "_");
assert(TypeParameters.size() == 8 &&
"Wrong number of literals in SPIR-V builtin image type");
StringRef SampledType = TypeParameters[0];
unsigned Dim, Depth, Arrayed, Multisampled, Sampled, Format, AccessQual;
bool AreParameterLiteralsValid =
!(TypeParameters[1].getAsInteger(10, Dim) ||
TypeParameters[2].getAsInteger(10, Depth) ||
TypeParameters[3].getAsInteger(10, Arrayed) ||
TypeParameters[4].getAsInteger(10, Multisampled) ||
TypeParameters[5].getAsInteger(10, Sampled) ||
TypeParameters[6].getAsInteger(10, Format) ||
TypeParameters[7].getAsInteger(10, AccessQual));
assert(AreParameterLiteralsValid &&
"Invalid format of SPIR-V image type parameter literals.");
return std::unique_ptr<SPIRV::ImageType>(new SPIRV::ImageType{
Name, SampledType, SPIRV::AccessQualifier::AccessQualifier(AccessQual),
SPIRV::Dim::Dim(Dim), static_cast<bool>(Arrayed),
static_cast<bool>(Depth), static_cast<bool>(Multisampled),
static_cast<bool>(Sampled), SPIRV::ImageFormat::ImageFormat(Format)});
}
static std::unique_ptr<const SPIRV::PipeType>
lookupOrParseBuiltinPipeType(StringRef Name) {
if (Name.startswith("opencl.")) {
// Lookup OpenCL builtin pipe type lowering details in TableGen records.
const SPIRV::PipeType *Record = SPIRV::lookupPipeType(Name);
return std::unique_ptr<SPIRV::PipeType>(new SPIRV::PipeType(*Record));
}
if (!Name.startswith("spirv."))
llvm_unreachable("Unknown builtin pipe type name/literal");
// Parse the access qualifier literal in the name of the SPIR-V pipe type.
// The name should have the following format:
// spirv.Pipe._AccessQualifier
// e.g. %spirv.Pipe._1
if (Name.endswith("_0"))
return std::unique_ptr<SPIRV::PipeType>(
new SPIRV::PipeType{Name, SPIRV::AccessQualifier::ReadOnly});
if (Name.endswith("_1"))
return std::unique_ptr<SPIRV::PipeType>(
new SPIRV::PipeType{Name, SPIRV::AccessQualifier::WriteOnly});
if (Name.endswith("_2"))
return std::unique_ptr<SPIRV::PipeType>(
new SPIRV::PipeType{Name, SPIRV::AccessQualifier::ReadWrite});
llvm_unreachable("Unknown pipe type access qualifier literal");
}
//===----------------------------------------------------------------------===//
// Implementation functions for builtin types.
//===----------------------------------------------------------------------===//
static SPIRVType *getNonParametrizedType(const StructType *OpaqueType,
const SPIRV::DemangledType *TypeRecord,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
unsigned Opcode = TypeRecord->Opcode;
// Create or get an existing type from GlobalRegistry.
return GR->getOrCreateOpTypeByOpcode(OpaqueType, MIRBuilder, Opcode);
}
static SPIRVType *getSamplerType(MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Create or get an existing type from GlobalRegistry.
return GR->getOrCreateOpTypeSampler(MIRBuilder);
}
static SPIRVType *getPipeType(const StructType *OpaqueType,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
// Lookup pipe type lowering details in TableGen records or parse the
// name/literal for details.
std::unique_ptr<const SPIRV::PipeType> Record =
lookupOrParseBuiltinPipeType(OpaqueType->getName());
// Create or get an existing type from GlobalRegistry.
return GR->getOrCreateOpTypePipe(MIRBuilder, Record.get()->Qualifier);
}
static SPIRVType *
getImageType(const StructType *OpaqueType,
SPIRV::AccessQualifier::AccessQualifier AccessQual,
MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
// Lookup image type lowering details in TableGen records or parse the
// name/literal for details.
std::unique_ptr<const SPIRV::ImageType> Record =
lookupOrParseBuiltinImageType(OpaqueType->getName());
SPIRVType *SampledType =
GR->getOrCreateSPIRVTypeByName(Record.get()->SampledType, MIRBuilder);
return GR->getOrCreateOpTypeImage(
MIRBuilder, SampledType, Record.get()->Dimensionality,
Record.get()->Depth, Record.get()->Arrayed, Record.get()->Multisampled,
Record.get()->Sampled, Record.get()->Format,
AccessQual == SPIRV::AccessQualifier::WriteOnly
? SPIRV::AccessQualifier::WriteOnly
: Record.get()->Qualifier);
}
static SPIRVType *getSampledImageType(const StructType *OpaqueType,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
StringRef TypeParametersString =
OpaqueType->getName().substr(strlen("spirv.SampledImage."));
LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
Type *ImageOpaqueType = StructType::getTypeByName(
Context, "spirv.Image." + TypeParametersString.str());
SPIRVType *TargetImageType =
GR->getOrCreateSPIRVType(ImageOpaqueType, MIRBuilder);
return GR->getOrCreateOpTypeSampledImage(TargetImageType, MIRBuilder);
}
namespace SPIRV {
SPIRVType *lowerBuiltinType(const StructType *OpaqueType,
SPIRV::AccessQualifier::AccessQualifier AccessQual,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
assert(OpaqueType->hasName() &&
"Structs representing builtin types must have a parsable name");
unsigned NumStartingVRegs = MIRBuilder.getMRI()->getNumVirtRegs();
const StringRef Name = OpaqueType->getName();
LLVM_DEBUG(dbgs() << "Lowering builtin type: " << Name << "\n");
// Lookup the demangled builtin type in the TableGen records.
const SPIRV::DemangledType *TypeRecord = findBuiltinType(Name);
if (!TypeRecord)
report_fatal_error("Missing TableGen record for builtin type: " + Name);
// "Lower" the BuiltinType into TargetType. The following get<...>Type methods
// use the implementation details from TableGen records to either create a new
// OpType<...> machine instruction or get an existing equivalent SPIRVType
// from GlobalRegistry.
SPIRVType *TargetType;
switch (TypeRecord->Opcode) {
case SPIRV::OpTypeImage:
TargetType = getImageType(OpaqueType, AccessQual, MIRBuilder, GR);
break;
case SPIRV::OpTypePipe:
TargetType = getPipeType(OpaqueType, MIRBuilder, GR);
break;
case SPIRV::OpTypeDeviceEvent:
TargetType = GR->getOrCreateOpTypeDeviceEvent(MIRBuilder);
break;
case SPIRV::OpTypeSampler:
TargetType = getSamplerType(MIRBuilder, GR);
break;
case SPIRV::OpTypeSampledImage:
TargetType = getSampledImageType(OpaqueType, MIRBuilder, GR);
break;
default:
TargetType = getNonParametrizedType(OpaqueType, TypeRecord, MIRBuilder, GR);
break;
}
// Emit OpName instruction if a new OpType<...> instruction was added
// (equivalent type was not found in GlobalRegistry).
if (NumStartingVRegs < MIRBuilder.getMRI()->getNumVirtRegs())
buildOpName(GR->getSPIRVTypeID(TargetType), OpaqueType->getName(),
MIRBuilder);
return TargetType;
}
} // namespace SPIRV
} // namespace llvm