MLIR
17.0.0git
|
Namespaces | |
detail | |
Classes | |
class | MemorySpaceToStorageClassConverter |
Type converter for converting numeric MemRef memory spaces into SPIR-V symbolic ones. More... | |
class | InterfaceVarABIAttr |
An attribute that specifies the information regarding the interface variable: descriptor set, binding, storage class. More... | |
class | VerCapExtAttr |
An attribute that specifies the SPIR-V (version, capabilities, extensions) triple. More... | |
class | TargetEnvAttr |
An attribute that specifies the target version, allowed extensions and capabilities, and resource limits. More... | |
class | SPIRVType |
class | ScalarType |
class | CompositeType |
class | ArrayType |
class | ImageType |
class | PointerType |
class | RuntimeArrayType |
class | SampledImageType |
class | StructType |
SPIR-V struct type. More... | |
class | CooperativeMatrixNVType |
class | JointMatrixINTELType |
class | MatrixType |
class | TargetEnv |
A wrapper class around a spirv::TargetEnvAttr to provide query methods for allowed version/capabilities/extensions. More... | |
struct | SerializationOptions |
struct | ElementwiseOpPattern |
Converts elementwise unary, binary and ternary standard operations to SPIR-V operations. More... | |
struct | BlockMergeInfo |
A struct for containing a header block's merge and continue targets. More... | |
struct | DebugLine |
A struct for containing OpLine instruction information. More... | |
struct | DeferredStructTypeInfo |
A "deferred struct type" is a struct type with one or more member types not known when the Deserializer first encounters the struct. More... | |
struct | SpecConstOperationMaterializationInfo |
A struct that collects the info needed to materialize/emit a SpecConstantOperation op. More... | |
class | Deserializer |
A SPIR-V module serializer. More... | |
class | Serializer |
A SPIR-V module serializer. More... | |
Typedefs | |
using | MemorySpaceToStorageClassMap = std::function< std::optional< spirv::StorageClass >(Attribute)> |
Mapping from numeric MemRef memory spaces into SPIR-V symbolic ones. More... | |
using | SymbolRenameListener = function_ref< void(spirv::ModuleOp originalModule, StringRef oldSymbol, StringRef newSymbol)> |
The listener function to receive symbol renaming events. More... | |
using | GetTargetEnvFn = std::function< spirv::TargetEnvAttr(spirv::ModuleOp)> |
Creates an operation pass that unifies access of multiple aliased resources into access of one single resource. More... | |
using | BlockMergeInfoMap = DenseMap< Block *, BlockMergeInfo > |
Map from a selection/loop's header block to its merge (and continue) target. More... | |
Functions | |
std::optional< spirv::StorageClass > | mapMemorySpaceToVulkanStorageClass (Attribute) |
Maps MemRef memory spaces to storage classes for Vulkan-flavored SPIR-V using the default rule. More... | |
std::optional< unsigned > | mapVulkanStorageClassToMemorySpace (spirv::StorageClass) |
Maps storage classes for Vulkan-flavored SPIR-V to MemRef memory spaces using the default rule. More... | |
std::optional< spirv::StorageClass > | mapMemorySpaceToOpenCLStorageClass (Attribute) |
Maps MemRef memory spaces to storage classes for OpenCL-flavored SPIR-V using the default rule. More... | |
std::optional< unsigned > | mapOpenCLStorageClassToMemorySpace (spirv::StorageClass) |
Maps storage classes for OpenCL-flavored SPIR-V to MemRef memory spaces using the default rule. More... | |
std::unique_ptr< ConversionTarget > | getMemorySpaceToStorageClassTarget (MLIRContext &) |
Creates the target that populates legality of ops with MemRef types. More... | |
void | populateMemorySpaceToStorageClassPatterns (MemorySpaceToStorageClassConverter &typeConverter, RewritePatternSet &patterns) |
Appends to a pattern list additional patterns for converting numeric MemRef memory spaces into SPIR-V symbolic ones. More... | |
ArrayRef< Extension > | getImpliedExtensions (Version version) |
Returns the implied extensions for the given version. More... | |
ArrayRef< Capability > | getDirectImpliedCapabilities (Capability cap) |
Returns the directly implied capabilities for the given capability. More... | |
SmallVector< Capability, 0 > | getRecursiveImpliedCapabilities (Capability cap) |
Returns the recursively implied capabilities for the given capability. More... | |
void | populateSPIRVGLCanonicalizationPatterns (RewritePatternSet &results) |
Populates patterns to run canoncalization that involves GL ops. More... | |
llvm::hash_code | hash_value (const StructType::MemberDecorationInfo &memberDecorationInfo) |
StringRef | getInterfaceVarABIAttrName () |
Returns the attribute name for specifying argument ABI information. More... | |
InterfaceVarABIAttr | getInterfaceVarABIAttr (unsigned descriptorSet, unsigned binding, std::optional< StorageClass > storageClass, MLIRContext *context) |
Gets the InterfaceVarABIAttr given its fields. More... | |
bool | needsInterfaceVarABIAttrs (TargetEnvAttr targetAttr) |
Returns whether the given SPIR-V target (described by TargetEnvAttr) needs ABI attributes for interface variables (spirv.interface_var_abi). More... | |
StringRef | getEntryPointABIAttrName () |
Returns the attribute name for specifying entry point information. More... | |
EntryPointABIAttr | getEntryPointABIAttr (MLIRContext *context, ArrayRef< int32_t > workgroupSize={}, std::optional< int > subgroupSize={}) |
Gets the EntryPointABIAttr given its fields. More... | |
EntryPointABIAttr | lookupEntryPointABI (Operation *op) |
Queries the entry point ABI on the nearest function-like op containing the given op . More... | |
DenseI32ArrayAttr | lookupLocalWorkGroupSize (Operation *op) |
Queries the local workgroup size from entry point ABI on the nearest function-like op containing the given op . More... | |
ResourceLimitsAttr | getDefaultResourceLimits (MLIRContext *context) |
Returns a default resource limits attribute that uses numbers from "Table 46. Required Limits" of the Vulkan spec. More... | |
StringRef | getTargetEnvAttrName () |
Returns the attribute name for specifying SPIR-V target environment. More... | |
TargetEnvAttr | getDefaultTargetEnv (MLIRContext *context) |
Returns the default target environment: SPIR-V 1.0 with Shader capability and no extra extensions. More... | |
TargetEnvAttr | lookupTargetEnv (Operation *op) |
Queries the target environment recursively from enclosing symbol table ops containing the given op . More... | |
TargetEnvAttr | lookupTargetEnvOrDefault (Operation *op) |
Queries the target environment recursively from enclosing symbol table ops containing the given op or returns the default target environment as returned by getDefaultTargetEnv() if not provided. More... | |
AddressingModel | getAddressingModel (TargetEnvAttr targetAttr) |
Returns addressing model selected based on target environment. More... | |
FailureOr< ExecutionModel > | getExecutionModel (TargetEnvAttr targetAttr) |
Returns execution model selected based on target environment. More... | |
FailureOr< MemoryModel > | getMemoryModel (TargetEnvAttr targetAttr) |
Returns memory model selected based on target environment. More... | |
OwningOpRef< spirv::ModuleOp > | combine (ArrayRef< spirv::ModuleOp > inputModules, OpBuilder &combinedModuleBuilder, SymbolRenameListener symRenameListener) |
Combines a list of SPIR-V inputModules into one. More... | |
std::unique_ptr< OperationPass< spirv::ModuleOp > > | createUnifyAliasedResourcePass (GetTargetEnvFn getTargetEnv=nullptr) |
Value | getBuiltinVariableValue (Operation *op, BuiltIn builtin, Type integerType, OpBuilder &builder) |
Returns the value for the given builtin variable. More... | |
Value | getPushConstantValue (Operation *op, unsigned elementCount, unsigned offset, Type integerType, OpBuilder &builder) |
Gets the value at the given offset of the push constant storage with a total of elementCount integerType integers. More... | |
Value | linearizeIndex (ValueRange indices, ArrayRef< int64_t > strides, int64_t offset, Type integerType, Location loc, OpBuilder &builder) |
Generates IR to perform index linearization with the given indices and their corresponding strides , adding an initial offset . More... | |
Value | getElementPtr (SPIRVTypeConverter &typeConverter, MemRefType baseType, Value basePtr, ValueRange indices, Location loc, OpBuilder &builder) |
Performs the index computation to get to the element at indices of the memory pointed to by basePtr , using the layout map of baseType . More... | |
Value | getOpenCLElementPtr (SPIRVTypeConverter &typeConverter, MemRefType baseType, Value basePtr, ValueRange indices, Location loc, OpBuilder &builder) |
Value | getVulkanElementPtr (SPIRVTypeConverter &typeConverter, MemRefType baseType, Value basePtr, ValueRange indices, Location loc, OpBuilder &builder) |
void | populateSPIRVExpandExtendedMultiplicationPatterns (RewritePatternSet &patterns) |
Appends to a pattern list additional patterns to expand extended multiplication ops into regular arithmetic ops. More... | |
OwningOpRef< spirv::ModuleOp > | deserialize (ArrayRef< uint32_t > binary, MLIRContext *context) |
Deserializes the given SPIR-V binary module and creates a MLIR ModuleOp in the given context . More... | |
LogicalResult | serialize (ModuleOp module, SmallVectorImpl< uint32_t > &binary, const SerializationOptions &options={}) |
Serializes the given SPIR-V module and writes to binary . More... | |
void | appendModuleHeader (SmallVectorImpl< uint32_t > &header, spirv::Version version, uint32_t idBound) |
Appends a SPRI-V module header to header with the given version and idBound . More... | |
uint32_t | getPrefixedOpcode (uint32_t wordCount, spirv::Opcode opcode) |
Returns the word-count-prefixed opcode for an SPIR-V instruction. More... | |
void | encodeStringLiteralInto (SmallVectorImpl< uint32_t > &binary, StringRef literal) |
Encodes an SPIR-V literal string into the given binary vector. More... | |
StringRef | decodeStringLiteral (ArrayRef< uint32_t > words, unsigned &wordIndex) |
Decodes a string literal in words starting at wordIndex . More... | |
template<> | |
LogicalResult | Deserializer::processOp< spirv::EntryPointOp > (ArrayRef< uint32_t > words) |
template<> | |
LogicalResult | Deserializer::processOp< spirv::ExecutionModeOp > (ArrayRef< uint32_t > words) |
template<> | |
LogicalResult | Deserializer::processOp< spirv::FunctionCallOp > (ArrayRef< uint32_t > operands) |
template<> | |
LogicalResult | Deserializer::processOp< spirv::CopyMemoryOp > (ArrayRef< uint32_t > words) |
template<> | |
LogicalResult | Deserializer::processOp< spirv::GenericCastToPtrExplicitOp > (ArrayRef< uint32_t > words) |
template<> | |
LogicalResult | Serializer::processOp< spirv::EntryPointOp > (spirv::EntryPointOp op) |
template<> | |
LogicalResult | Serializer::processOp< spirv::ExecutionModeOp > (spirv::ExecutionModeOp op) |
template<> | |
LogicalResult | Serializer::processOp< spirv::FunctionCallOp > (spirv::FunctionCallOp op) |
template<> | |
LogicalResult | Serializer::processOp< spirv::CopyMemoryOp > (spirv::CopyMemoryOp op) |
template<> | |
LogicalResult | Serializer::processOp< spirv::GenericCastToPtrExplicitOp > (spirv::GenericCastToPtrExplicitOp op) |
void | encodeInstructionInto (SmallVectorImpl< uint32_t > &binary, spirv::Opcode op, ArrayRef< uint32_t > operands) |
Encodes an SPIR-V instruction with the given opcode and operands into the given binary vector. More... | |
template<> | |
LogicalResult | Serializer::processTypeDecoration< spirv::ArrayType > (Location loc, spirv::ArrayType type, uint32_t resultID) |
template<> | |
LogicalResult | Serializer::processTypeDecoration< spirv::RuntimeArrayType > (Location loc, spirv::RuntimeArrayType type, uint32_t resultID) |
Variables | |
constexpr unsigned | kHeaderWordCount = 5 |
SPIR-V binary header word count. More... | |
constexpr uint32_t | kMagicNumber = 0x07230203 |
SPIR-V magic number. More... | |
constexpr uint32_t | kGeneratorNumber = 22 |
The serializer tool ID registered to the Khronos Group. More... | |
using mlir::spirv::BlockMergeInfoMap = typedef DenseMap<Block *, BlockMergeInfo> |
Map from a selection/loop's header block to its merge (and continue) target.
Definition at line 60 of file Deserializer.h.
using mlir::spirv::GetTargetEnvFn = typedef std::function<spirv::TargetEnvAttr(spirv::ModuleOp)> |
using mlir::spirv::MemorySpaceToStorageClassMap = typedef std::function<std::optional<spirv::StorageClass>(Attribute)> |
Mapping from numeric MemRef memory spaces into SPIR-V symbolic ones.
Definition at line 25 of file MemRefToSPIRV.h.
using mlir::spirv::SymbolRenameListener = typedef function_ref<void( spirv::ModuleOp originalModule, StringRef oldSymbol, StringRef newSymbol)> |
The listener function to receive symbol renaming events.
originalModule
is the input spirv::ModuleOp that contains the renamed symbol. oldSymbol
and newSymbol
are the original and renamed symbol. Note that it's the responsibility of the caller to properly retain the storage underlying the passed StringRefs if the listener callback outlives this function call.
Definition at line 32 of file ModuleCombiner.h.
void mlir::spirv::appendModuleHeader | ( | SmallVectorImpl< uint32_t > & | header, |
spirv::Version | version, | ||
uint32_t | idBound | ||
) |
Appends a SPRI-V module header to header
with the given version
and idBound
.
Definition at line 18 of file SPIRVBinaryUtils.cpp.
References kGeneratorNumber, kMagicNumber, and MIN_VERSION_CASE.
Referenced by mlir::spirv::Serializer::collect().
OwningOpRef< spirv::ModuleOp > mlir::spirv::combine | ( | ArrayRef< spirv::ModuleOp > | inputModules, |
OpBuilder & | combinedModuleBuilder, | ||
SymbolRenameListener | symRenameListener | ||
) |
Combines a list of SPIR-V inputModules
into one.
Returns the combined module on success; returns a null module otherwise.
inputModules | the list of modules to combine. They won't be modified. |
combinedMdouleBuilder | an OpBuilder for building the combined module. |
symbRenameListener | a listener that gets called everytime a symbol in one of the input modules is renamed. |
To combine multiple SPIR-V modules, we move all the module-level ops from all the input modules into one big combined module. To that end, the combination process proceeds in 2 phases:
For the conflict resolution phase, the following rules are employed to resolve such conflicts:
For deduplication, the following 3 cases are taken into consideration:
In all cases, the references to the updated symbol (whether renamed or deduplicated) are also updated to reflect the change.
Definition at line 90 of file ModuleCombiner.cpp.
References computeHash(), mlir::OpBuilder::create(), mlir::failed(), mlir::OpBuilder::insert(), mlir::SymbolTable::replaceAllSymbolUses(), mlir::OpBuilder::setInsertionPointToStart(), and updateSymbolAndAllUses().
std::unique_ptr< mlir::OperationPass< spirv::ModuleOp > > mlir::spirv::createUnifyAliasedResourcePass | ( | spirv::GetTargetEnvFn | getTargetEnv = nullptr | ) |
Definition at line 614 of file UnifyAliasedResourcePass.cpp.
|
inline |
Decodes a string literal in words
starting at wordIndex
.
Update the latter to point to the position in words after the string literal.
Definition at line 47 of file SPIRVBinaryUtils.h.
Referenced by Deserializer::processOp< spirv::EntryPointOp >().
OwningOpRef< spirv::ModuleOp > mlir::spirv::deserialize | ( | ArrayRef< uint32_t > | binary, |
MLIRContext * | context | ||
) |
Deserializes the given SPIR-V binary
module and creates a MLIR ModuleOp in the given context
.
Returns the ModuleOp on success; otherwise, reports errors to the error handler registered with context
and returns a null module.
Definition at line 15 of file Deserialization.cpp.
References mlir::spirv::Deserializer::collect(), mlir::spirv::Deserializer::deserialize(), and mlir::failed().
Referenced by deserializeModule(), and roundTripModule().
LogicalResult mlir::spirv::Deserializer::processOp< spirv::CopyMemoryOp > | ( | ArrayRef< uint32_t > | words | ) |
Definition at line 455 of file DeserializeOps.cpp.
References mlir::emitError(), and mlir::success().
LogicalResult mlir::spirv::Deserializer::processOp< spirv::EntryPointOp > | ( | ArrayRef< uint32_t > | words | ) |
Definition at line 331 of file DeserializeOps.cpp.
References decodeStringLiteral(), mlir::emitError(), and mlir::success().
LogicalResult mlir::spirv::Deserializer::processOp< spirv::ExecutionModeOp > | ( | ArrayRef< uint32_t > | words | ) |
Definition at line 380 of file DeserializeOps.cpp.
References mlir::emitError(), and mlir::success().
LogicalResult mlir::spirv::Deserializer::processOp< spirv::FunctionCallOp > | ( | ArrayRef< uint32_t > | operands | ) |
Definition at line 413 of file DeserializeOps.cpp.
References mlir::emitError(), and mlir::success().
LogicalResult mlir::spirv::Deserializer::processOp< spirv::GenericCastToPtrExplicitOp > | ( | ArrayRef< uint32_t > | words | ) |
Definition at line 527 of file DeserializeOps.cpp.
References mlir::Operation::create(), mlir::emitError(), mlir::Operation::getResult(), and mlir::success().
void mlir::spirv::encodeInstructionInto | ( | SmallVectorImpl< uint32_t > & | binary, |
spirv::Opcode | op, | ||
ArrayRef< uint32_t > | operands | ||
) |
Encodes an SPIR-V instruction with the given opcode
and operands
into the given binary
vector.
Definition at line 78 of file Serializer.cpp.
References getPrefixedOpcode().
Referenced by Serializer::processOp< spirv::CopyMemoryOp >(), Serializer::processOp< spirv::EntryPointOp >(), Serializer::processOp< spirv::ExecutionModeOp >(), Serializer::processOp< spirv::FunctionCallOp >(), and Serializer::processOp< spirv::GenericCastToPtrExplicitOp >().
void mlir::spirv::encodeStringLiteralInto | ( | SmallVectorImpl< uint32_t > & | binary, |
StringRef | literal | ||
) |
Encodes an SPIR-V literal
string into the given binary
vector.
Definition at line 66 of file SPIRVBinaryUtils.cpp.
Referenced by Serializer::processOp< spirv::EntryPointOp >().
spirv::AddressingModel mlir::spirv::getAddressingModel | ( | spirv::TargetEnvAttr | targetAttr | ) |
Returns addressing model selected based on target environment.
Definition at line 208 of file TargetAndABI.cpp.
References mlir::spirv::TargetEnvAttr::getCapabilities().
Value mlir::spirv::getBuiltinVariableValue | ( | Operation * | op, |
BuiltIn | builtin, | ||
Type | integerType, | ||
OpBuilder & | builder | ||
) |
Returns the value for the given builtin
variable.
This function gets or inserts the global variable associated for the builtin within the nearest symbol table enclosing op
. Returns null Value on error.
spirv::ResourceLimitsAttr mlir::spirv::getDefaultResourceLimits | ( | MLIRContext * | context | ) |
Returns a default resource limits attribute that uses numbers from "Table 46. Required Limits" of the Vulkan spec.
Definition at line 157 of file TargetAndABI.cpp.
References mlir::Builder::getI32ArrayAttr().
Referenced by getDefaultTargetEnv().
spirv::TargetEnvAttr mlir::spirv::getDefaultTargetEnv | ( | MLIRContext * | context | ) |
Returns the default target environment: SPIR-V 1.0 with Shader capability and no extra extensions.
Definition at line 174 of file TargetAndABI.cpp.
References mlir::spirv::TargetEnvAttr::get(), mlir::spirv::VerCapExtAttr::get(), getDefaultResourceLimits(), and mlir::spirv::TargetEnvAttr::kUnknownDeviceID.
Referenced by lookupTargetEnvOrDefault().
ArrayRef<Capability> mlir::spirv::getDirectImpliedCapabilities | ( | Capability | cap | ) |
Returns the directly implied capabilities for the given capability.
These capabilities are implicitly declared by the given capability.
Value mlir::spirv::getElementPtr | ( | SPIRVTypeConverter & | typeConverter, |
MemRefType | baseType, | ||
Value | basePtr, | ||
ValueRange | indices, | ||
Location | loc, | ||
OpBuilder & | builder | ||
) |
Performs the index computation to get to the element at indices
of the memory pointed to by basePtr
, using the layout map of baseType
.
Returns null if index computation cannot be performed.
Definition at line 812 of file SPIRVConversion.cpp.
References mlir::SPIRVTypeConverter::allows(), getOpenCLElementPtr(), and getVulkanElementPtr().
spirv::EntryPointABIAttr mlir::spirv::getEntryPointABIAttr | ( | MLIRContext * | context, |
ArrayRef< int32_t > | workgroupSize = {} , |
||
std::optional< int > | subgroupSize = {} |
||
) |
Gets the EntryPointABIAttr given its fields.
Definition at line 124 of file TargetAndABI.cpp.
References mlir::detail::DenseArrayAttrImpl< int32_t >::get().
StringRef mlir::spirv::getEntryPointABIAttrName | ( | ) |
Returns the attribute name for specifying entry point information.
Definition at line 121 of file TargetAndABI.cpp.
Referenced by lookupEntryPointABI(), lowerAsEntryFunction(), and lowerEntryPointABIAttr().
FailureOr< spirv::ExecutionModel > mlir::spirv::getExecutionModel | ( | spirv::TargetEnvAttr | targetAttr | ) |
Returns execution model selected based on target environment.
Returns failure if it cannot be selected.
Definition at line 225 of file TargetAndABI.cpp.
References mlir::failure(), and mlir::spirv::TargetEnvAttr::getCapabilities().
Referenced by lowerEntryPointABIAttr().
ArrayRef<Extension> mlir::spirv::getImpliedExtensions | ( | Version | version | ) |
Returns the implied extensions for the given version.
These extensions are incorporated into the current version so they are implicitly declared when targeting the given version.
Referenced by mlir::spirv::TargetEnv::TargetEnv().
InterfaceVarABIAttr mlir::spirv::getInterfaceVarABIAttr | ( | unsigned | descriptorSet, |
unsigned | binding, | ||
std::optional< StorageClass > | storageClass, | ||
MLIRContext * | context | ||
) |
Gets the InterfaceVarABIAttr given its fields.
Referenced by getDefaultABIAttrs().
StringRef mlir::spirv::getInterfaceVarABIAttrName | ( | ) |
Returns the attribute name for specifying argument ABI information.
Definition at line 99 of file TargetAndABI.cpp.
Referenced by getDefaultABIAttrs(), lowerAsEntryFunction(), and verifyRegionAttribute().
FailureOr< spirv::MemoryModel > mlir::spirv::getMemoryModel | ( | spirv::TargetEnvAttr | targetAttr | ) |
Returns memory model selected based on target environment.
Returns failure if it cannot be selected.
Definition at line 236 of file TargetAndABI.cpp.
References mlir::failure(), mlir::spirv::TargetEnvAttr::getCapabilities(), and mlir::gpu::amd::OpenCL.
std::unique_ptr< ConversionTarget > mlir::spirv::getMemorySpaceToStorageClassTarget | ( | MLIRContext & | context | ) |
Creates the target that populates legality of ops with MemRef types.
Definition at line 240 of file MapMemRefStorageClassPass.cpp.
References isLegalOp().
Value mlir::spirv::getOpenCLElementPtr | ( | SPIRVTypeConverter & | typeConverter, |
MemRefType | baseType, | ||
Value | basePtr, | ||
ValueRange | indices, | ||
Location | loc, | ||
OpBuilder & | builder | ||
) |
Definition at line 777 of file SPIRVConversion.cpp.
References mlir::Type::cast(), mlir::OpBuilder::create(), mlir::failed(), mlir::SPIRVTypeConverter::getIndexType(), mlir::getStridesAndOffset(), mlir::Value::getType(), getZero(), mlir::Type::isa(), and linearizeIndex().
Referenced by getElementPtr().
uint32_t mlir::spirv::getPrefixedOpcode | ( | uint32_t | wordCount, |
spirv::Opcode | opcode | ||
) |
Returns the word-count-prefixed opcode for an SPIR-V instruction.
Definition at line 61 of file SPIRVBinaryUtils.cpp.
Referenced by encodeInstructionInto().
Value mlir::spirv::getPushConstantValue | ( | Operation * | op, |
unsigned | elementCount, | ||
unsigned | offset, | ||
Type | integerType, | ||
OpBuilder & | builder | ||
) |
Gets the value at the given offset
of the push constant storage with a total of elementCount
integerType
integers.
A global variable will be created in the nearest symbol table enclosing op
for the push constant storage if not existing. Load ops will be created via the given builder
to load values from the push constant. Returns null Value on error.
Definition at line 696 of file SPIRVConversion.cpp.
References mlir::OpBuilder::create(), mlir::Operation::emitError(), mlir::Region::front(), mlir::Builder::getI32IntegerAttr(), mlir::Operation::getLoc(), mlir::SymbolTable::getNearestSymbolTable(), getOrInsertPushConstantVariable(), mlir::Operation::getParentOp(), mlir::Operation::getRegion(), and getZero().
SmallVector<Capability, 0> mlir::spirv::getRecursiveImpliedCapabilities | ( | Capability | cap | ) |
Returns the recursively implied capabilities for the given capability.
These capabilities are implicitly declared by the given capability. Compared to the above function, this function collects implied capabilities recursively: if an implicitly declared capability implicitly declares a third one, the third one will also be returned.
Referenced by mlir::spirv::TargetEnv::TargetEnv().
StringRef mlir::spirv::getTargetEnvAttrName | ( | ) |
Returns the attribute name for specifying SPIR-V target environment.
Definition at line 172 of file TargetAndABI.cpp.
Referenced by lookupTargetEnv().
Value mlir::spirv::getVulkanElementPtr | ( | SPIRVTypeConverter & | typeConverter, |
MemRefType | baseType, | ||
Value | basePtr, | ||
ValueRange | indices, | ||
Location | loc, | ||
OpBuilder & | builder | ||
) |
Definition at line 746 of file SPIRVConversion.cpp.
References mlir::OpBuilder::create(), mlir::failed(), mlir::SPIRVTypeConverter::getIndexType(), mlir::getStridesAndOffset(), getZero(), and linearizeIndex().
Referenced by getElementPtr().
llvm::hash_code mlir::spirv::hash_value | ( | const StructType::MemberDecorationInfo & | memberDecorationInfo | ) |
Definition at line 1144 of file SPIRVTypes.cpp.
References mlir::spirv::StructType::MemberDecorationInfo::decoration, and mlir::spirv::StructType::MemberDecorationInfo::memberIndex.
Value mlir::spirv::linearizeIndex | ( | ValueRange | indices, |
ArrayRef< int64_t > | strides, | ||
int64_t | offset, | ||
Type | integerType, | ||
Location | loc, | ||
OpBuilder & | builder | ||
) |
Generates IR to perform index linearization with the given indices
and their corresponding strides
, adding an initial offset
.
Definition at line 722 of file SPIRVConversion.cpp.
References mlir::OpBuilder::create(), and mlir::detail::enumerate().
Referenced by getOpenCLElementPtr(), and getVulkanElementPtr().
spirv::EntryPointABIAttr mlir::spirv::lookupEntryPointABI | ( | Operation * | op | ) |
Queries the entry point ABI on the nearest function-like op containing the given op
.
Returns null attribute if not found.
Definition at line 136 of file TargetAndABI.cpp.
References mlir::Operation::getAttrOfType(), getEntryPointABIAttrName(), and mlir::Operation::getParentOp().
Referenced by lookupLocalWorkGroupSize().
DenseI32ArrayAttr mlir::spirv::lookupLocalWorkGroupSize | ( | Operation * | op | ) |
Queries the local workgroup size from entry point ABI on the nearest function-like op containing the given op
.
Returns null attribute if not found.
Definition at line 149 of file TargetAndABI.cpp.
References lookupEntryPointABI().
spirv::TargetEnvAttr mlir::spirv::lookupTargetEnv | ( | Operation * | op | ) |
Queries the target environment recursively from enclosing symbol table ops containing the given op
.
Definition at line 184 of file TargetAndABI.cpp.
References mlir::Operation::getAttrOfType(), mlir::SymbolTable::getNearestSymbolTable(), mlir::Operation::getParentOp(), and getTargetEnvAttrName().
Referenced by lookupTargetEnvOrDefault(), and lowerEntryPointABIAttr().
spirv::TargetEnvAttr mlir::spirv::lookupTargetEnvOrDefault | ( | Operation * | op | ) |
Queries the target environment recursively from enclosing symbol table ops containing the given op
or returns the default target environment as returned by getDefaultTargetEnv() if not provided.
Definition at line 200 of file TargetAndABI.cpp.
References mlir::Operation::getContext(), getDefaultTargetEnv(), and lookupTargetEnv().
Referenced by getDefaultABIAttrs().
std::optional< spirv::StorageClass > mlir::spirv::mapMemorySpaceToOpenCLStorageClass | ( | Attribute | memorySpaceAttr | ) |
Maps MemRef memory spaces to storage classes for OpenCL-flavored SPIR-V using the default rule.
Returns std::nullopt if the memory space is unknown.
Definition at line 114 of file MapMemRefStorageClassPass.cpp.
References mlir::Attribute::dyn_cast(), OPENCL_STORAGE_SPACE_MAP_LIST, and STORAGE_SPACE_MAP_FN.
std::optional< spirv::StorageClass > mlir::spirv::mapMemorySpaceToVulkanStorageClass | ( | Attribute | memorySpaceAttr | ) |
Maps MemRef memory spaces to storage classes for Vulkan-flavored SPIR-V using the default rule.
Returns std::nullopt if the memory space is unknown.
Definition at line 60 of file MapMemRefStorageClassPass.cpp.
References mlir::Attribute::dyn_cast(), STORAGE_SPACE_MAP_FN, and VULKAN_STORAGE_SPACE_MAP_LIST.
std::optional< unsigned > mlir::spirv::mapOpenCLStorageClassToMemorySpace | ( | spirv::StorageClass | storageClass | ) |
Maps storage classes for OpenCL-flavored SPIR-V to MemRef memory spaces using the default rule.
Returns std::nullopt if the storage class is unsupported.
Definition at line 141 of file MapMemRefStorageClassPass.cpp.
References OPENCL_STORAGE_SPACE_MAP_LIST, and STORAGE_SPACE_MAP_FN.
std::optional< unsigned > mlir::spirv::mapVulkanStorageClassToMemorySpace | ( | spirv::StorageClass | storageClass | ) |
Maps storage classes for Vulkan-flavored SPIR-V to MemRef memory spaces using the default rule.
Returns std::nullopt if the storage class is unsupported.
Definition at line 87 of file MapMemRefStorageClassPass.cpp.
References STORAGE_SPACE_MAP_FN, and VULKAN_STORAGE_SPACE_MAP_LIST.
bool mlir::spirv::needsInterfaceVarABIAttrs | ( | spirv::TargetEnvAttr | targetAttr | ) |
Returns whether the given SPIR-V target (described by TargetEnvAttr) needs ABI attributes for interface variables (spirv.interface_var_abi).
Definition at line 111 of file TargetAndABI.cpp.
References mlir::spirv::TargetEnvAttr::getCapabilities().
Referenced by getDefaultABIAttrs().
void mlir::spirv::populateMemorySpaceToStorageClassPatterns | ( | spirv::MemorySpaceToStorageClassConverter & | typeConverter, |
RewritePatternSet & | patterns | ||
) |
Appends to a pattern list additional patterns for converting numeric MemRef memory spaces into SPIR-V symbolic ones.
Definition at line 297 of file MapMemRefStorageClassPass.cpp.
References mlir::RewritePatternSet::add(), and mlir::RewritePatternSet::getContext().
void mlir::spirv::populateSPIRVExpandExtendedMultiplicationPatterns | ( | RewritePatternSet & | patterns | ) |
Appends to a pattern list additional patterns to expand extended multiplication ops into regular arithmetic ops.
Extended multiplication ops are not supported by the WebGPU Shading Language (WGSL).
Definition at line 190 of file SPIRVWebGPUTransforms.cpp.
References mlir::RewritePatternSet::add(), and mlir::RewritePatternSet::getContext().
void mlir::spirv::populateSPIRVGLCanonicalizationPatterns | ( | RewritePatternSet & | results | ) |
Populates patterns to run canoncalization that involves GL ops.
These patterns cannot be run in default canonicalization because GL ops aren't always available. So they should be involed specifically when needed.
Definition at line 25 of file SPIRVGLCanonicalization.cpp.
References mlir::RewritePatternSet::add(), and mlir::RewritePatternSet::getContext().
LogicalResult mlir::spirv::serialize | ( | ModuleOp | module, |
SmallVectorImpl< uint32_t > & | binary, | ||
const SerializationOptions & | options = {} |
||
) |
Serializes the given SPIR-V module
and writes to binary
.
On failure, reports errors to the error handler registered with the MLIR context for module
.
Referenced by roundTripModule(), and serializeModule().
LogicalResult mlir::spirv::Serializer::processOp< spirv::CopyMemoryOp > | ( | spirv::CopyMemoryOp | op | ) |
Definition at line 628 of file SerializeOps.cpp.
References encodeInstructionInto(), mlir::failed(), mlir::failure(), and mlir::success().
LogicalResult mlir::spirv::Serializer::processOp< spirv::EntryPointOp > | ( | spirv::EntryPointOp | op | ) |
Definition at line 537 of file SerializeOps.cpp.
References encodeInstructionInto(), encodeStringLiteralInto(), mlir::FlatSymbolRefAttr::getValue(), and mlir::success().
LogicalResult mlir::spirv::Serializer::processOp< spirv::ExecutionModeOp > | ( | spirv::ExecutionModeOp | op | ) |
Definition at line 572 of file SerializeOps.cpp.
References encodeInstructionInto(), and mlir::success().
LogicalResult mlir::spirv::Serializer::processOp< spirv::FunctionCallOp > | ( | spirv::FunctionCallOp | op | ) |
Definition at line 601 of file SerializeOps.cpp.
References encodeInstructionInto(), mlir::failed(), mlir::failure(), mlir::Type::isa(), and mlir::success().
LogicalResult mlir::spirv::Serializer::processOp< spirv::GenericCastToPtrExplicitOp > | ( | spirv::GenericCastToPtrExplicitOp | op | ) |
Definition at line 673 of file SerializeOps.cpp.
References mlir::Type::cast(), encodeInstructionInto(), mlir::failed(), mlir::failure(), and mlir::success().
LogicalResult mlir::spirv::Serializer::processTypeDecoration< spirv::ArrayType > | ( | Location | loc, |
spirv::ArrayType | type, | ||
uint32_t | resultID | ||
) |
Definition at line 270 of file Serializer.cpp.
References mlir::spirv::ArrayType::getArrayStride(), and mlir::success().
LogicalResult mlir::spirv::Serializer::processTypeDecoration< spirv::RuntimeArrayType > | ( | Location | loc, |
spirv::RuntimeArrayType | type, | ||
uint32_t | resultID | ||
) |
Definition at line 280 of file Serializer.cpp.
References mlir::spirv::RuntimeArrayType::getArrayStride(), and mlir::success().
|
constexpr |
The serializer tool ID registered to the Khronos Group.
Definition at line 31 of file SPIRVBinaryUtils.h.
Referenced by appendModuleHeader().
|
constexpr |
SPIR-V binary header word count.
Definition at line 25 of file SPIRVBinaryUtils.h.
Referenced by mlir::spirv::Serializer::collect().
|
constexpr |
SPIR-V magic number.
Definition at line 28 of file SPIRVBinaryUtils.h.
Referenced by appendModuleHeader().