Namespaces
	detail

Classes
class	ArrayType

struct	BlockMergeInfo
	A struct for containing a header block's merge and continue targets. More...

class	CompositeType

class	CooperativeMatrixNVType

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...

class	Deserializer
	A SPIR-V module serializer. More...

class	ElementwiseOpPattern
	Converts elementwise unary, binary and ternary standard operations to SPIR-V operations. More...

class	ImageType

class	InterfaceVarABIAttr
	An attribute that specifies the information regarding the interface variable: descriptor set, binding, storage class. More...

class	MatrixType

class	PointerType

class	RuntimeArrayType

class	SampledImageType

class	ScalarType

struct	SerializationOptions

class	Serializer
	A SPIR-V module serializer. More...

struct	SpecConstOperationMaterializationInfo
	A struct that collects the info needed to materialize/emit a SpecConstantOperation op. More...

class	SPIRVType

class	StructType
	SPIR-V struct type. More...

class	TargetEnv
	A wrapper class around a spirv::TargetEnvAttr to provide query methods for allowed version/capabilities/extensions. More...

class	TargetEnvAttr
	An attribute that specifies the target version, allowed extensions and capabilities, and resource limits. More...

class	VerCapExtAttr
	An attribute that specifies the SPIR-V (version, capabilities, extensions) triple. More...

Typedefs
using	SymbolRenameListener = function_ref< void(spirv::ModuleOp originalModule, StringRef oldSymbol, StringRef newSymbol)>
	The listener function to receive symbol renaming events. More...

using	BlockMergeInfoMap = DenseMap< Block *, BlockMergeInfo >
	Map from a selection/loop's header block to its merge (and continue) target. More...

Functions
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	populateSPIRVGLSLCanonicalizationPatterns (RewritePatternSet &results)
	Populates patterns to run canoncalization that involves GLSL 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, 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 (spv.interface_var_abi). More...

StringRef	getEntryPointABIAttrName ()
	Returns the attribute name for specifying entry point information. More...

EntryPointABIAttr	getEntryPointABIAttr (ArrayRef< int32_t > localSize, MLIRContext *context)
	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...

DenseIntElementsAttr	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<> >	createCanonicalizeGLSLPass ()
	Creates a pass to run canoncalization patterns that involve GLSL ops. More...

std::unique_ptr< OperationPass< mlir::ModuleOp > >	createDecorateSPIRVCompositeTypeLayoutPass ()
	Creates a module pass that converts composite types used by objects in the StorageBuffer, PhysicalStorageBuffer, Uniform, and PushConstant storage classes with layout information. More...

std::unique_ptr< OperationPass< spirv::ModuleOp > >	createUpdateVersionCapabilityExtensionPass ()
	Creates an operation pass that deduces and attaches the minimal version/ capabilities/extensions requirements for spv.module ops. More...

std::unique_ptr< OperationPass< spirv::ModuleOp > >	createLowerABIAttributesPass ()
	Creates an operation pass that lowers the ABI attributes specified during SPIR-V Lowering. More...

std::unique_ptr< OperationPass< spirv::ModuleOp > >	createRewriteInsertsPass ()
	Creates an operation pass that rewrites sequential chains of spv.CompositeInsert into spv.CompositeConstruct. More...

std::unique_ptr< OperationPass< spirv::ModuleOp > >	createUnifyAliasedResourcePass ()
	Creates an operation pass that unifies access of multiple aliased resources into access of one single resource. More...

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...

spirv::AccessChainOp	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...

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::ControlBarrierOp > (ArrayRef< uint32_t > operands)

template<>
LogicalResult	Deserializer::processOp< spirv::FunctionCallOp > (ArrayRef< uint32_t > operands)

template<>
LogicalResult	Deserializer::processOp< spirv::MemoryBarrierOp > (ArrayRef< uint32_t > operands)

template<>
LogicalResult	Deserializer::processOp< spirv::CopyMemoryOp > (ArrayRef< uint32_t > words)

template<>
LogicalResult	Serializer::processOp< spirv::EntryPointOp > (spirv::EntryPointOp op)

template<>
LogicalResult	Serializer::processOp< spirv::ControlBarrierOp > (spirv::ControlBarrierOp op)

template<>
LogicalResult	Serializer::processOp< spirv::ExecutionModeOp > (spirv::ExecutionModeOp op)

template<>
LogicalResult	Serializer::processOp< spirv::MemoryBarrierOp > (spirv::MemoryBarrierOp op)

template<>
LogicalResult	Serializer::processOp< spirv::FunctionCallOp > (spirv::FunctionCallOp op)

template<>
LogicalResult	Serializer::processOp< spirv::CopyMemoryOp > (spirv::CopyMemoryOp 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...

Typedef Documentation

◆ BlockMergeInfoMap

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 59 of file Deserializer.h.

◆ SymbolRenameListener

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 33 of file ModuleCombiner.h.

Function Documentation

◆ appendModuleHeader()

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().

◆ combine()

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.

Parameters

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:

resolve conflicts between pairs of ops from different modules,
deduplicate equivalent ops/sub-ops in the merged module.

For the conflict resolution phase, the following rules are employed to resolve such conflicts:

If 2 spv.func's have the same symbol name, then rename one of the functions.
If an spv.func and another op have the same symbol name, then rename the other symbol.
If none of the 2 conflicting ops are spv.func, then rename either.

For deduplication, the following 3 cases are taken into consideration:

If 2 spv.GlobalVariable's have either the same descriptor set + binding or the same build_in attribute value, then replace one of them using the other.
If 2 spv.SpecConstant's have the same spec_id attribute value, then replace one of them using the other.
Deduplicating functions are not supported right now.

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().

Referenced by combineOneSpec().

◆ createCanonicalizeGLSLPass()

std::unique_ptr< OperationPass<> > mlir::spirv::createCanonicalizeGLSLPass ( )

Creates a pass to run canoncalization patterns that involve GLSL ops.

These patterns cannot be run in default canonicalization because GLSL ops aren't always available. So they should be involed specifically when needed.

Definition at line 32 of file CanonicalizeGLSLPass.cpp.

◆ createDecorateSPIRVCompositeTypeLayoutPass()

std::unique_ptr< OperationPass< ModuleOp > > mlir::spirv::createDecorateSPIRVCompositeTypeLayoutPass ( )

Creates a module pass that converts composite types used by objects in the StorageBuffer, PhysicalStorageBuffer, Uniform, and PushConstant storage classes with layout information.

Right now this pass only supports Vulkan layout rules.

Definition at line 143 of file DecorateCompositeTypeLayoutPass.cpp.

◆ createLowerABIAttributesPass()

std::unique_ptr< OperationPass< spirv::ModuleOp > > mlir::spirv::createLowerABIAttributesPass ( )

Creates an operation pass that lowers the ABI attributes specified during SPIR-V Lowering.

Specifically,

Creates the global variables for arguments of entry point function using the specification in the spv.interface_var_abi attribute for each argument.
Inserts the EntryPointOp and the ExecutionModeOp for entry point functions using the specification in the spv.entry_point_abi attribute.

Definition at line 293 of file LowerABIAttributesPass.cpp.

◆ createRewriteInsertsPass()

std::unique_ptr< mlir::OperationPass< spirv::ModuleOp > > mlir::spirv::createRewriteInsertsPass ( )

Creates an operation pass that rewrites sequential chains of spv.CompositeInsert into spv.CompositeConstruct.

Definition at line 113 of file RewriteInsertsPass.cpp.

◆ createUnifyAliasedResourcePass()

std::unique_ptr< mlir::OperationPass< spirv::ModuleOp > > mlir::spirv::createUnifyAliasedResourcePass ( )

Creates an operation pass that unifies access of multiple aliased resources into access of one single resource.

Definition at line 452 of file UnifyAliasedResourcePass.cpp.

◆ createUpdateVersionCapabilityExtensionPass()

std::unique_ptr< OperationPass< spirv::ModuleOp > > mlir::spirv::createUpdateVersionCapabilityExtensionPass ( )

Creates an operation pass that deduces and attaches the minimal version/ capabilities/extensions requirements for spv.module ops.

For each spv.module op, this pass requires a spv.target_env attribute on it or an enclosing module-like op to drive the deduction. The reason is that an op can be enabled by multiple extensions/capabilities. So we need to know which one to pick. spv.target_env gives the hard limit as for what the target environment can support; this pass deduces what are actually needed for a specific spv.module op.

Definition at line 181 of file UpdateVCEPass.cpp.

◆ decodeStringLiteral()

StringRef mlir::spirv::decodeStringLiteral	(	ArrayRef< uint32_t >	words,
		unsigned &	wordIndex
	)

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 mlir::spirv::Deserializer::collect(), and Deserializer::processOp< spirv::EntryPointOp >().

◆ deserialize()

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().

◆ Deserializer::processOp< spirv::ControlBarrierOp >()

template<>

LogicalResult mlir::spirv::Deserializer::processOp< spirv::ControlBarrierOp > ( ArrayRef< uint32_t > operands )

Definition at line 405 of file DeserializeOps.cpp.

References mlir::emitError(), and mlir::success().

◆ Deserializer::processOp< spirv::CopyMemoryOp >()

template<>

LogicalResult mlir::spirv::Deserializer::processOp< spirv::CopyMemoryOp > ( ArrayRef< uint32_t > words )

Definition at line 501 of file DeserializeOps.cpp.

References mlir::emitError(), and mlir::success().

◆ Deserializer::processOp< spirv::EntryPointOp >()

template<>

LogicalResult mlir::spirv::Deserializer::processOp< spirv::EntryPointOp > ( ArrayRef< uint32_t > words )

Definition at line 329 of file DeserializeOps.cpp.

References decodeStringLiteral(), mlir::emitError(), and mlir::success().

◆ Deserializer::processOp< spirv::ExecutionModeOp >()

template<>

LogicalResult mlir::spirv::Deserializer::processOp< spirv::ExecutionModeOp > ( ArrayRef< uint32_t > words )

Definition at line 372 of file DeserializeOps.cpp.

References mlir::emitError(), and mlir::success().

◆ Deserializer::processOp< spirv::FunctionCallOp >()

template<>

LogicalResult mlir::spirv::Deserializer::processOp< spirv::FunctionCallOp > ( ArrayRef< uint32_t > operands )

Definition at line 434 of file DeserializeOps.cpp.

References mlir::emitError(), mlir::success(), and value.

◆ Deserializer::processOp< spirv::MemoryBarrierOp >()

template<>

LogicalResult mlir::spirv::Deserializer::processOp< spirv::MemoryBarrierOp > ( ArrayRef< uint32_t > operands )

Definition at line 476 of file DeserializeOps.cpp.

References mlir::emitError(), and mlir::success().

◆ encodeInstructionInto()

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 75 of file Serializer.cpp.

References getPrefixedOpcode().

Referenced by mlir::spirv::Serializer::printValueIDMap(), Serializer::processOp< spirv::ControlBarrierOp >(), Serializer::processOp< spirv::CopyMemoryOp >(), Serializer::processOp< spirv::EntryPointOp >(), Serializer::processOp< spirv::ExecutionModeOp >(), Serializer::processOp< spirv::FunctionCallOp >(), Serializer::processOp< spirv::MemoryBarrierOp >(), Serializer::processTypeDecoration< spirv::RuntimeArrayType >(), and visitInPrettyBlockOrder().

◆ encodeStringLiteralInto()

void mlir::spirv::encodeStringLiteralInto	(	SmallVectorImpl< uint32_t > &	binary,
		StringRef	literal
	)

Encodes an SPIR-V literal string into the given binary vector.

Definition at line 65 of file SPIRVBinaryUtils.cpp.

Referenced by mlir::spirv::Serializer::printValueIDMap(), and Serializer::processOp< spirv::EntryPointOp >().

◆ getAddressingModel()

spirv::AddressingModel mlir::spirv::getAddressingModel ( spirv::TargetEnvAttr targetAttr )

Returns addressing model selected based on target environment.

Definition at line 202 of file TargetAndABI.cpp.

References mlir::spirv::TargetEnvAttr::getCapabilities().

Referenced by getDefaultABIAttrs(), and mlir::spirv::TargetEnv::operator TargetEnvAttr().

◆ getBuiltinVariableValue()

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.

Referenced by getLocalInvocationDimSize(), and getOrInsertBuiltinVariable().

◆ getDefaultResourceLimits()

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 155 of file TargetAndABI.cpp.

Referenced by getDefaultTargetEnv(), and mlir::spirv::TargetEnv::operator TargetEnvAttr().

◆ 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 168 of file TargetAndABI.cpp.

References mlir::spirv::VerCapExtAttr::get(), mlir::spirv::TargetEnvAttr::get(), getDefaultResourceLimits(), and mlir::spirv::TargetEnvAttr::kUnknownDeviceID.

Referenced by lookupTargetEnvOrDefault(), mlir::spirv::TargetEnv::operator TargetEnvAttr(), and verifyConstantType().

◆ getDirectImpliedCapabilities()

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.

◆ getElementPtr()

spirv::AccessChainOp 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 784 of file SPIRVConversion.cpp.

References mlir::OpBuilder::create(), mlir::failed(), mlir::SPIRVTypeConverter::getIndexType(), mlir::getStridesAndOffset(), and linearizeIndex().

Referenced by castBoolToIntN(), and getLocalInvocationDimSize().

◆ getEntryPointABIAttr()

spirv::EntryPointABIAttr mlir::spirv::getEntryPointABIAttr	(	ArrayRef< int32_t >	localSize,
		MLIRContext *	context
	)

Gets the EntryPointABIAttr given its fields.

Definition at line 122 of file TargetAndABI.cpp.

Referenced by mlir::spirv::TargetEnv::operator TargetEnvAttr().

◆ getEntryPointABIAttrName()

StringRef mlir::spirv::getEntryPointABIAttrName ( )

Returns the attribute name for specifying entry point information.

Definition at line 119 of file TargetAndABI.cpp.

Referenced by lookupEntryPointABI(), lowerAsEntryFunction(), lowerEntryPointABIAttr(), mlir::spirv::TargetEnv::operator TargetEnvAttr(), and print().

◆ getExecutionModel()

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 214 of file TargetAndABI.cpp.

References mlir::failure(), and mlir::spirv::TargetEnvAttr::getCapabilities().

Referenced by lowerEntryPointABIAttr(), and mlir::spirv::TargetEnv::operator TargetEnvAttr().

◆ getImpliedExtensions()

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().

◆ getInterfaceVarABIAttr()

InterfaceVarABIAttr mlir::spirv::getInterfaceVarABIAttr	(	unsigned	descriptorSet,
		unsigned	binding,
		Optional< StorageClass >	storageClass,
		MLIRContext *	context
	)

Gets the InterfaceVarABIAttr given its fields.

Referenced by getDefaultABIAttrs(), getInterfaceVarABIAttrName(), and mlir::spirv::TargetEnv::operator TargetEnvAttr().

◆ getInterfaceVarABIAttrName()

StringRef mlir::spirv::getInterfaceVarABIAttrName ( )

Returns the attribute name for specifying argument ABI information.

Definition at line 97 of file TargetAndABI.cpp.

References mlir::spirv::InterfaceVarABIAttr::get(), and getInterfaceVarABIAttr().

Referenced by getDefaultABIAttrs(), lowerAsEntryFunction(), lowerEntryPointABIAttr(), mlir::spirv::TargetEnv::operator TargetEnvAttr(), and verifyRegionAttribute().

◆ getMemoryModel()

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 225 of file TargetAndABI.cpp.

References mlir::failure(), mlir::spirv::TargetEnvAttr::getCapabilities(), and mlir::gpu::amd::OpenCL.

Referenced by getDefaultABIAttrs(), and mlir::spirv::TargetEnv::operator TargetEnvAttr().

◆ getPrefixedOpcode()

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 60 of file SPIRVBinaryUtils.cpp.

Referenced by encodeInstructionInto().

◆ getPushConstantValue()

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 734 of file SPIRVConversion.cpp.

References mlir::Operation::emitError(), mlir::Region::front(), mlir::Operation::getLoc(), mlir::SymbolTable::getNearestSymbolTable(), getOrInsertPushConstantVariable(), mlir::Operation::getParentOp(), and mlir::Operation::getRegion().

◆ getRecursiveImpliedCapabilities()

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().

◆ getTargetEnvAttrName()

StringRef mlir::spirv::getTargetEnvAttrName ( )

Returns the attribute name for specifying SPIR-V target environment.

Definition at line 166 of file TargetAndABI.cpp.

Referenced by lookupTargetEnv(), mlir::spirv::TargetEnv::operator TargetEnvAttr(), and print().

◆ hash_value()

llvm::hash_code mlir::spirv::hash_value ( const StructType::MemberDecorationInfo & memberDecorationInfo )

Definition at line 1058 of file SPIRVTypes.cpp.

References mlir::spirv::StructType::MemberDecorationInfo::decoration, and mlir::spirv::StructType::MemberDecorationInfo::memberIndex.

Referenced by mlir::python::PyIntegerSet::createFromCapsule(), llvm::DenseMapInfo< mlir::tblgen::DagNode >::getHashValue(), llvm::DenseMapInfo< mlir::tblgen::DagLeaf >::getHashValue(), mlir::detail::DenseIntOrFPElementsAttrStorage::getKey(), mlir::detail::DenseStringElementsAttrStorage::getKey(), mlir::detail::DenseIntOrFPElementsAttrStorage::getKeyForBoolData(), mlir::detail::IntegerTypeStorage::hashKey(), mlir::detail::DynamicTypeStorage::hashKey(), mlir::detail::DynamicAttrStorage::hashKey(), mlir::pdl_to_pdl_interp::OperandGroupPosition::hashKey(), mlir::pdl_to_pdl_interp::OperationPosition::hashKey(), mlir::pdl_to_pdl_interp::ResultGroupPosition::hashKey(), mlir::pdl_to_pdl_interp::UsersPosition::hashKey(), mlirTypeIDHashValue(), and mlir::python::PySymbolTable::walkSymbolTables().

◆ linearizeIndex()

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 760 of file SPIRVConversion.cpp.

References mlir::OpBuilder::create(), and mlir::detail::enumerate().

Referenced by getElementPtr().

◆ lookupEntryPointABI()

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 134 of file TargetAndABI.cpp.

References mlir::Operation::getAttrOfType(), getEntryPointABIAttrName(), and mlir::Operation::getParentOp().

Referenced by getDefaultABIAttrs(), lookupLocalWorkGroupSize(), and mlir::spirv::TargetEnv::operator TargetEnvAttr().

◆ lookupLocalWorkGroupSize()

DenseIntElementsAttr 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 147 of file TargetAndABI.cpp.

References lookupEntryPointABI().

Referenced by getLocalInvocationDimSize(), and mlir::spirv::TargetEnv::operator TargetEnvAttr().

◆ lookupTargetEnv()

spirv::TargetEnvAttr mlir::spirv::lookupTargetEnv ( Operation * op )

Queries the target environment recursively from enclosing symbol table ops containing the given op.

Definition at line 178 of file TargetAndABI.cpp.

References mlir::Operation::getAttrOfType(), mlir::SymbolTable::getNearestSymbolTable(), mlir::Operation::getParentOp(), and getTargetEnvAttrName().

Referenced by checkAndUpdateCapabilityRequirements(), lookupTargetEnvOrDefault(), lowerEntryPointABIAttr(), and mlir::spirv::TargetEnv::operator TargetEnvAttr().

◆ lookupTargetEnvOrDefault()

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 194 of file TargetAndABI.cpp.

References mlir::Operation::getContext(), getDefaultTargetEnv(), and lookupTargetEnv().

Referenced by getDefaultABIAttrs(), mlir::spirv::TargetEnv::operator TargetEnvAttr(), mlir::arith::populateArithmeticToSPIRVPatterns(), and verifyConstantType().

◆ needsInterfaceVarABIAttrs()

bool mlir::spirv::needsInterfaceVarABIAttrs ( spirv::TargetEnvAttr targetAttr )

Returns whether the given SPIR-V target (described by TargetEnvAttr) needs ABI attributes for interface variables (spv.interface_var_abi).

Definition at line 109 of file TargetAndABI.cpp.

References mlir::spirv::TargetEnvAttr::getCapabilities().

Referenced by getDefaultABIAttrs(), and mlir::spirv::TargetEnv::operator TargetEnvAttr().

◆ populateSPIRVGLSLCanonicalizationPatterns()

void mlir::spirv::populateSPIRVGLSLCanonicalizationPatterns ( RewritePatternSet & results )

Populates patterns to run canoncalization that involves GLSL ops.

These patterns cannot be run in default canonicalization because GLSL ops aren't always available. So they should be involed specifically when needed.

Definition at line 25 of file SPIRVGLSLCanonicalization.cpp.

References mlir::RewritePatternSet::add(), and mlir::RewritePatternSet::getContext().

◆ serialize()

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().

◆ Serializer::processOp< spirv::ControlBarrierOp >()

template<>

LogicalResult mlir::spirv::Serializer::processOp< spirv::ControlBarrierOp > ( spirv::ControlBarrierOp op )

Definition at line 570 of file SerializeOps.cpp.

References encodeInstructionInto(), mlir::failure(), and mlir::success().

◆ Serializer::processOp< spirv::CopyMemoryOp >()

template<>

LogicalResult mlir::spirv::Serializer::processOp< spirv::CopyMemoryOp > ( spirv::CopyMemoryOp op )

Definition at line 666 of file SerializeOps.cpp.

References encodeInstructionInto(), mlir::failed(), mlir::failure(), and mlir::success().

◆ Serializer::processOp< spirv::EntryPointOp >()

template<>

LogicalResult mlir::spirv::Serializer::processOp< spirv::EntryPointOp > ( spirv::EntryPointOp op )

Definition at line 536 of file SerializeOps.cpp.

References encodeInstructionInto(), encodeStringLiteralInto(), mlir::FlatSymbolRefAttr::getValue(), and mlir::success().

◆ Serializer::processOp< spirv::ExecutionModeOp >()

template<>

LogicalResult mlir::spirv::Serializer::processOp< spirv::ExecutionModeOp > ( spirv::ExecutionModeOp op )

Definition at line 591 of file SerializeOps.cpp.

References encodeInstructionInto(), and mlir::success().

◆ Serializer::processOp< spirv::FunctionCallOp >()

template<>

LogicalResult mlir::spirv::Serializer::processOp< spirv::FunctionCallOp > ( spirv::FunctionCallOp op )

Definition at line 639 of file SerializeOps.cpp.

References encodeInstructionInto(), mlir::failed(), mlir::failure(), mlir::Type::isa(), mlir::success(), and value.

◆ Serializer::processOp< spirv::MemoryBarrierOp >()

template<>

LogicalResult mlir::spirv::Serializer::processOp< spirv::MemoryBarrierOp > ( spirv::MemoryBarrierOp op )

Definition at line 620 of file SerializeOps.cpp.

References encodeInstructionInto(), mlir::failure(), and mlir::success().

◆ Serializer::processTypeDecoration< spirv::ArrayType >()

template<>

LogicalResult mlir::spirv::Serializer::processTypeDecoration< spirv::ArrayType >	(	Location	loc,
		spirv::ArrayType	type,
		uint32_t	resultID
	)

Definition at line 263 of file Serializer.cpp.

References mlir::success().

◆ Serializer::processTypeDecoration< spirv::RuntimeArrayType >()

template<>

LogicalResult mlir::spirv::Serializer::processTypeDecoration< spirv::RuntimeArrayType >	(	Location	loc,
		spirv::RuntimeArrayType	type,
		uint32_t	resultID
	)

Variable Documentation

◆ kGeneratorNumber

constexpr uint32_t mlir::spirv::kGeneratorNumber = 22

The serializer tool ID registered to the Khronos Group.

Definition at line 31 of file SPIRVBinaryUtils.h.

Referenced by appendModuleHeader().

◆ kHeaderWordCount

constexpr unsigned mlir::spirv::kHeaderWordCount = 5

SPIR-V binary header word count.

Definition at line 25 of file SPIRVBinaryUtils.h.

Referenced by mlir::spirv::Serializer::collect(), and mlir::spirv::Deserializer::collect().

◆ kMagicNumber

constexpr uint32_t mlir::spirv::kMagicNumber = 0x07230203

SPIR-V magic number.

Definition at line 28 of file SPIRVBinaryUtils.h.

Referenced by appendModuleHeader(), and mlir::spirv::Deserializer::collect().

Namespaces

Classes

Typedefs

Functions

Variables

Typedef Documentation

◆ BlockMergeInfoMap

◆ SymbolRenameListener

Function Documentation

◆ appendModuleHeader()

◆ combine()

◆ createCanonicalizeGLSLPass()

◆ createDecorateSPIRVCompositeTypeLayoutPass()

◆ createLowerABIAttributesPass()

◆ createRewriteInsertsPass()

◆ createUnifyAliasedResourcePass()

◆ createUpdateVersionCapabilityExtensionPass()

◆ decodeStringLiteral()

◆ deserialize()

◆ Deserializer::processOp< spirv::ControlBarrierOp >()

◆ Deserializer::processOp< spirv::CopyMemoryOp >()

◆ Deserializer::processOp< spirv::EntryPointOp >()

◆ Deserializer::processOp< spirv::ExecutionModeOp >()

◆ Deserializer::processOp< spirv::FunctionCallOp >()

◆ Deserializer::processOp< spirv::MemoryBarrierOp >()

◆ encodeInstructionInto()

◆ encodeStringLiteralInto()

◆ getAddressingModel()

◆ getBuiltinVariableValue()

◆ getDefaultResourceLimits()

◆ getDefaultTargetEnv()

◆ getDirectImpliedCapabilities()

◆ getElementPtr()

◆ getEntryPointABIAttr()

◆ getEntryPointABIAttrName()

◆ getExecutionModel()

◆ getImpliedExtensions()

◆ getInterfaceVarABIAttr()

◆ getInterfaceVarABIAttrName()

◆ getMemoryModel()

◆ getPrefixedOpcode()

◆ getPushConstantValue()

◆ getRecursiveImpliedCapabilities()

◆ getTargetEnvAttrName()

◆ hash_value()

◆ linearizeIndex()

◆ lookupEntryPointABI()

◆ lookupLocalWorkGroupSize()

◆ lookupTargetEnv()

◆ lookupTargetEnvOrDefault()

◆ needsInterfaceVarABIAttrs()

◆ populateSPIRVGLSLCanonicalizationPatterns()

◆ serialize()

◆ Serializer::processOp< spirv::ControlBarrierOp >()

◆ Serializer::processOp< spirv::CopyMemoryOp >()

◆ Serializer::processOp< spirv::EntryPointOp >()

◆ Serializer::processOp< spirv::ExecutionModeOp >()

◆ Serializer::processOp< spirv::FunctionCallOp >()

◆ Serializer::processOp< spirv::MemoryBarrierOp >()

◆ Serializer::processTypeDecoration< spirv::ArrayType >()

◆ Serializer::processTypeDecoration< spirv::RuntimeArrayType >()

Variable Documentation

◆ kGeneratorNumber

◆ kHeaderWordCount

◆ kMagicNumber