Namespaces
	AttrNames

	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	CooperativeMatrixType

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	convertMemRefTypesAndAttrs (Operation *op, MemorySpaceToStorageClassConverter &typeConverter)
	Converts all MemRef types and attributes in the op, as decided by the `typeConverter`. 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={}, std::optional< int > targetWidth={})
	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, bool use64bitAddress)
	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, StringRef prefix="__builtin__", StringRef suffix="__")
	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 (const 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 (const SPIRVTypeConverter &typeConverter, MemRefType baseType, Value basePtr, ValueRange indices, Location loc, OpBuilder &builder)

Value	getVulkanElementPtr (const SPIRVTypeConverter &typeConverter, MemRefType baseType, Value basePtr, ValueRange indices, Location loc, OpBuilder &builder)

int	getComputeVectorSize (int64_t size)

SmallVector< int64_t >	getNativeVectorShapeImpl (vector::ReductionOp op)

SmallVector< int64_t >	getNativeVectorShapeImpl (vector::TransposeOp op)

std::optional< SmallVector< int64_t > >	getNativeVectorShape (Operation *op)

LogicalResult	unrollVectorsInSignatures (Operation *op)

LogicalResult	unrollVectorsInFuncBodies (Operation *op)

void	populateSPIRVExpandExtendedMultiplicationPatterns (RewritePatternSet &patterns)
	Appends patterns to expand extended multiplication and adition ops into regular arithmetic ops. More...

void	populateSPIRVExpandNonFiniteArithmeticPatterns (RewritePatternSet &patterns)
	Appends patterns to expand non-finite arithmetic ops `IsNan` and `IsInf`. 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...

void	registerSPIRVTargetInterfaceExternalModels (DialectRegistry &registry)
	Registers the `TargetAttrInterface` for the `#spirv.target_env` attribute in the given registry. More...

void	registerSPIRVTargetInterfaceExternalModels (MLIRContext &context)
	Registers the `TargetAttrInterface` for the `#spirv.target_env` attribute in the registry associated with the given context. More...

template<typename T >
static StringRef	stringifyTypeName ()

template<>
StringRef	stringifyTypeName< IntegerType > ()

template<>
StringRef	stringifyTypeName< FloatType > ()

template<typename AtomicOpTy , typename ExpectedElementType >
static LogicalResult	verifyAtomicUpdateOp (Operation *op)

static LogicalResult	verifyCastOp (Operation *op, bool requireSameBitWidth=true, bool skipBitWidthCheck=false)

template<typename EnumAttrClass , typename EnumClass >
static ParseResult	parseControlAttribute (OpAsmParser &parser, OperationState &state, StringRef attrName=spirv::attributeName< EnumClass >())
	Parses Function, Selection and Loop control attributes. More...

static bool	hasOneBranchOpTo (Block &srcBlock, Block &dstBlock)
	Returns true if the given `srcBlock` contains only one `spirv.Branch` to the given `dstBlock`. More...

static bool	isMergeBlock (Block &block)
	Returns true if the given `block` only contains one `spirv.mlir.merge` op. More...

static bool	hasOtherMerge (Region &region)
	Returns true if a `spirv.mlir.merge` op outside the merge block. More...

static LogicalResult	verifyCoopMatrixAccess (Operation *op, Type pointer, Type coopMatrix, spirv::MemoryAccessAttr memoryOperand)

template<typename OpTy >
static LogicalResult	verifyGroupNonUniformArithmeticOp (Operation *groupOp)

template<typename OpTy >
static LogicalResult	verifyGroupNonUniformShuffleOp (OpTy op)

template<typename Op >
static LogicalResult	verifyGroupOp (Op op)

template<typename IntegerDotProductOpTy >
static LogicalResult	verifyIntegerDotProduct (Operation *op)

static std::optional< spirv::Version >	getIntegerDotProductMinVersion ()

static std::optional< spirv::Version >	getIntegerDotProductMaxVersion ()

static SmallVector< ArrayRef< spirv::Extension >, 1 >	getIntegerDotProductExtensions ()

template<typename IntegerDotProductOpTy >
static SmallVector< ArrayRef< spirv::Capability >, 1 >	getIntegerDotProductCapabilities (Operation *op)

template<typename MemoryOpTy >
ParseResult	parseMemoryAccessAttributes (OpAsmParser &parser, OperationState &state)
	Parses optional memory access (a.k.a. More...

template<typename MemoryOpTy >
static ParseResult	parseSourceMemoryAccessAttributes (OpAsmParser &parser, OperationState &state)

template<typename MemoryOpTy >
static void	printSourceMemoryAccessAttribute (MemoryOpTy memoryOp, OpAsmPrinter &printer, SmallVectorImpl< StringRef > &elidedAttrs, std::optional< spirv::MemoryAccess > memoryAccessAtrrValue=std::nullopt, std::optional< uint32_t > alignmentAttrValue=std::nullopt)

template<typename MemoryOpTy >
static void	printMemoryAccessAttribute (MemoryOpTy memoryOp, OpAsmPrinter &printer, SmallVectorImpl< StringRef > &elidedAttrs, std::optional< spirv::MemoryAccess > memoryAccessAtrrValue=std::nullopt, std::optional< uint32_t > alignmentAttrValue=std::nullopt)

template<typename LoadStoreOpTy >
static LogicalResult	verifyLoadStorePtrAndValTypes (LoadStoreOpTy op, Value ptr, Value val)

template<typename MemoryOpTy >
static LogicalResult	verifyMemoryAccessAttribute (MemoryOpTy memoryOp)

template<typename MemoryOpTy >
static LogicalResult	verifySourceMemoryAccessAttribute (MemoryOpTy memoryOp)

static Type	getElementPtrType (Type type, ValueRange indices, Location baseLoc)

template<typename Op >
static void	printAccessChain (Op op, ValueRange indices, OpAsmPrinter &printer)

template<typename Op >
static LogicalResult	verifyAccessChain (Op accessChainOp, ValueRange indices)

static bool	isNestedInFunctionOpInterface (Operation *op)
	Returns true if the given op is a function-like op or nested in a function-like op without a module-like op in the middle. More...

static bool	isDirectInModuleLikeOp (Operation *op)
	Returns true if the given op is an module-like op that maintains a symbol table. More...

static Type	getUnaryOpResultType (Type operandType)
	Result of a logical op must be a scalar or vector of boolean type. More...

static ParseResult	parseImageOperands (OpAsmParser &parser, spirv::ImageOperandsAttr &attr)

static void	printImageOperands (OpAsmPrinter &printer, Operation *imageOp, spirv::ImageOperandsAttr attr)

unsigned	getBitWidth (Type type)
	Returns the bit width of the `type`. More...

void	printVariableDecorations (Operation *op, OpAsmPrinter &printer, SmallVectorImpl< StringRef > &elidedAttrs)

LogicalResult	extractValueFromConstOp (Operation *op, int32_t &value)

LogicalResult	verifyMemorySemantics (Operation *op, spirv::MemorySemantics memorySemantics)

ParseResult	parseVariableDecorations (OpAsmParser &parser, OperationState &state)

template<typename Ty >
ArrayAttr	getStrArrayAttrForEnumList (Builder &builder, ArrayRef< Ty > enumValues, function_ref< StringRef(Ty)> stringifyFn)

template<typename EnumClass , typename ParserType >
ParseResult	parseEnumKeywordAttr (EnumClass &value, ParserType &parser, StringRef attrName=spirv::attributeName< EnumClass >())
	Parses the next keyword in `parser` as an enumerant of the given `EnumClass`. More...

template<typename EnumClass >
ParseResult	parseEnumStrAttr (EnumClass &value, OpAsmParser &parser, StringRef attrName=spirv::attributeName< EnumClass >())
	Parses the next string attribute in `parser` as an enumerant of the given `EnumClass`. More...

template<typename EnumAttrClass , typename EnumClass = typename EnumAttrClass::ValueType>
ParseResult	parseEnumStrAttr (EnumClass &value, OpAsmParser &parser, OperationState &state, StringRef attrName=spirv::attributeName< EnumClass >())
	Parses the next string attribute in `parser` as an enumerant of the given `EnumClass` and inserts the enumerant into `state` as an 32-bit integer attribute with the enum class's name as attribute name. More...

template<typename EnumAttrClass , typename EnumClass = typename EnumAttrClass::ValueType>
ParseResult	parseEnumKeywordAttr (EnumClass &value, OpAsmParser &parser, OperationState &state, StringRef attrName=spirv::attributeName< EnumClass >())
	Parses the next keyword in `parser` as an enumerant of the given `EnumClass` and inserts the enumerant into `state` as an 32-bit integer attribute with the enum class's name as attribute name. 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...

static std::string	getDecorationName (StringRef attrName)

template<typename AttrTy , typename EmitF >
LogicalResult	processDecorationList (Location loc, Decoration decoration, Attribute attrList, StringRef attrName, EmitF emitter)

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

◆ GetTargetEnvFn

using mlir::spirv::GetTargetEnvFn = typedef std::function<spirv::TargetEnvAttr(spirv::ModuleOp)>

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

Definition at line 36 of file Passes.h.

◆ MemorySpaceToStorageClassMap

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.

◆ 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 32 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 19 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 spirv.func's have the same symbol name, then rename one of the functions.
If an spirv.func and another op have the same symbol name, then rename the other symbol.
If none of the 2 conflicting ops are spirv.func, then rename either.

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

If 2 spirv.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 spirv.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::OpBuilder::insert(), mlir::SymbolTable::replaceAllSymbolUses(), mlir::OpBuilder::setInsertionPointToStart(), and updateSymbolAndAllUses().

◆ convertMemRefTypesAndAttrs()

void mlir::spirv::convertMemRefTypesAndAttrs	(	Operation *	op,
		MemorySpaceToStorageClassConverter &	typeConverter
	)

Converts all MemRef types and attributes in the op, as decided by the typeConverter.

Definition at line 248 of file MapMemRefStorageClassPass.cpp.

References mlir::detail::AttrTypeReplacerBase< Concrete >::addReplacement(), mlir::TypeConverter::convertType(), and mlir::detail::AttrTypeReplacerBase< Concrete >::recursivelyReplaceElementsIn().

◆ createUnifyAliasedResourcePass()

std::unique_ptr< mlir::OperationPass< spirv::ModuleOp > > mlir::spirv::createUnifyAliasedResourcePass ( spirv::GetTargetEnvFn getTargetEnv = nullptr )

Definition at line 627 of file UnifyAliasedResourcePass.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 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(), and mlir::spirv::Deserializer::deserialize().

Referenced by deserializeModule(), and roundTripModule().

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

template<>

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

Definition at line 453 of file DeserializeOps.cpp.

References mlir::emitError().

◆ 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::get().

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

template<>

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

Definition at line 378 of file DeserializeOps.cpp.

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

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

template<>

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

Definition at line 411 of file DeserializeOps.cpp.

References mlir::emitError(), mlir::get(), and mlir::getType().

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

template<>

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

Definition at line 526 of file DeserializeOps.cpp.

References mlir::Operation::create(), mlir::emitError(), mlir::Operation::getResult(), and mlir::getType().

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

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

Referenced by Serializer::processOp< spirv::EntryPointOp >().

◆ extractValueFromConstOp()

LogicalResult mlir::spirv::extractValueFromConstOp	(	Operation *	op,
		int32_t &	value
	)

Definition at line 50 of file SPIRVOps.cpp.

Referenced by getElementPtrType(), and verifyGroupNonUniformArithmeticOp().

◆ getAddressingModel()

spirv::AddressingModel mlir::spirv::getAddressingModel	(	spirv::TargetEnvAttr	targetAttr,
		bool	use64bitAddress
	)

Returns addressing model selected based on target environment.

Definition at line 205 of file TargetAndABI.cpp.

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

◆ getBitWidth()

unsigned mlir::spirv::getBitWidth ( Type type )

inline

Returns the bit width of the type.

Definition at line 14 of file SPIRVOpUtils.h.

References mlir::Type::getIntOrFloatBitWidth(), and mlir::Type::isIntOrFloat().

Referenced by verifyIntegerDotProduct().

◆ getBuiltinVariableValue()

Value mlir::spirv::getBuiltinVariableValue	(	Operation *	op,
		BuiltIn	builtin,
		Type	integerType,
		OpBuilder &	builder,
		StringRef	prefix = `"__builtin__"`,
		StringRef	suffix = `"__"`
	)

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.

The global name being generated will be mangled using preffix and suffix.

◆ getComputeVectorSize()

int mlir::spirv::getComputeVectorSize ( int64_t size )

Definition at line 1305 of file SPIRVConversion.cpp.

Referenced by getNativeVectorShape(), and getNativeVectorShapeImpl().

◆ getDecorationName()

static std::string mlir::spirv::getDecorationName ( StringRef attrName )

static

Definition at line 212 of file Serializer.cpp.

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

References mlir::get(), and mlir::Builder::getI32ArrayAttr().

Referenced by getDefaultTargetEnv().

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

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

Referenced by lookupTargetEnvOrDefault().

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

Value mlir::spirv::getElementPtr	(	const 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 1287 of file SPIRVConversion.cpp.

References mlir::SPIRVTypeConverter::allows(), getOpenCLElementPtr(), and getVulkanElementPtr().

◆ getElementPtrType()

static Type mlir::spirv::getElementPtrType	(	Type	type,
		ValueRange	indices,
		Location	baseLoc
	)

static

Definition at line 263 of file MemoryOps.cpp.

References mlir::emitError(), extractValueFromConstOp(), mlir::spirv::PointerType::get(), and mlir::Operation::getName().

Referenced by verifyAccessChain().

◆ getEntryPointABIAttr()

spirv::EntryPointABIAttr mlir::spirv::getEntryPointABIAttr	(	MLIRContext *	context,
		ArrayRef< int32_t >	workgroupSize = `{}`,
		std::optional< int >	subgroupSize = `{}`,
		std::optional< int >	targetWidth = `{}`
	)

Gets the EntryPointABIAttr given its fields.

targetWidth is used by several execution modes. It is the element width of floating-point operations. Refer to Execution Mode in SPIR-V specification. https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#_execution_mode

Definition at line 120 of file TargetAndABI.cpp.

References mlir::detail::DenseArrayAttrImpl< int32_t >::get(), mlir::get(), and subgroupSize.

◆ getEntryPointABIAttrName()

StringRef mlir::spirv::getEntryPointABIAttrName ( )

Returns the attribute name for specifying entry point information.

Definition at line 118 of file TargetAndABI.cpp.

Referenced by lookupEntryPointABI(), lowerAsEntryFunction(), and lowerEntryPointABIAttr().

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

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

Referenced by lowerEntryPointABIAttr().

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

◆ getIntegerDotProductCapabilities()

template<typename IntegerDotProductOpTy >

static SmallVector<ArrayRef<spirv::Capability>, 1> mlir::spirv::getIntegerDotProductCapabilities ( Operation * op )

static

Definition at line 92 of file IntegerDotProductOps.cpp.

References mlir::Operation::getAttr(), mlir::Operation::getName(), mlir::Operation::getOperand(), and mlir::Value::getType().

◆ getIntegerDotProductExtensions()

static SmallVector<ArrayRef<spirv::Extension>, 1> mlir::spirv::getIntegerDotProductExtensions ( )

static

Definition at line 83 of file IntegerDotProductOps.cpp.

◆ getIntegerDotProductMaxVersion()

static std::optional<spirv::Version> mlir::spirv::getIntegerDotProductMaxVersion ( )

static

Definition at line 78 of file IntegerDotProductOps.cpp.

◆ getIntegerDotProductMinVersion()

static std::optional<spirv::Version> mlir::spirv::getIntegerDotProductMinVersion ( )

static

Definition at line 74 of file IntegerDotProductOps.cpp.

◆ getInterfaceVarABIAttr()

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

Gets the InterfaceVarABIAttr given its fields.

Referenced by getDefaultABIAttrs().

◆ getInterfaceVarABIAttrName()

StringRef mlir::spirv::getInterfaceVarABIAttrName ( )

Returns the attribute name for specifying argument ABI information.

Definition at line 96 of file TargetAndABI.cpp.

Referenced by getDefaultABIAttrs(), lowerAsEntryFunction(), 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 233 of file TargetAndABI.cpp.

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

◆ getMemorySpaceToStorageClassTarget()

std::unique_ptr< ConversionTarget > mlir::spirv::getMemorySpaceToStorageClassTarget ( MLIRContext & context )

Creates the target that populates legality of ops with MemRef types.

Definition at line 242 of file MapMemRefStorageClassPass.cpp.

References isLegalOp().

◆ getNativeVectorShape()

std::optional< SmallVector< int64_t > > mlir::spirv::getNativeVectorShape ( Operation * op )

Definition at line 1332 of file SPIRVConversion.cpp.

References getComputeVectorSize(), getNativeVectorShapeImpl(), mlir::Operation::getNumResults(), mlir::Operation::getResultTypes(), and mlir::OpTrait::hasElementwiseMappableTraits().

Referenced by unrollVectorsInFuncBodies().

◆ getNativeVectorShapeImpl() [1/2]

SmallVector< int64_t > mlir::spirv::getNativeVectorShapeImpl ( vector::ReductionOp op )

Definition at line 1314 of file SPIRVConversion.cpp.

References getComputeVectorSize().

Referenced by getNativeVectorShape().

◆ getNativeVectorShapeImpl() [2/2]

SmallVector< int64_t > mlir::spirv::getNativeVectorShapeImpl ( vector::TransposeOp op )

Definition at line 1323 of file SPIRVConversion.cpp.

References getComputeVectorSize().

◆ getOpenCLElementPtr()

Value mlir::spirv::getOpenCLElementPtr	(	const SPIRVTypeConverter &	typeConverter,
		MemRefType	baseType,
		Value	basePtr,
		ValueRange	indices,
		Location	loc,
		OpBuilder &	builder
	)

Definition at line 1252 of file SPIRVConversion.cpp.

References mlir::OpBuilder::create(), mlir::SPIRVTypeConverter::getIndexType(), mlir::Value::getType(), getZero(), and mlir::affine::linearizeIndex().

Referenced by getElementPtr().

◆ 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 62 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 1170 of file SPIRVConversion.cpp.

References mlir::OpBuilder::create(), mlir::Operation::emitError(), mlir::Region::front(), mlir::Builder::getI32IntegerAttr(), mlir::Operation::getLoc(), mlir::SymbolTable::getNearestSymbolTable(), mlir::Operation::getParentOp(), mlir::Operation::getRegion(), and getZero().

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

◆ getStrArrayAttrForEnumList()

template<typename Ty >

ArrayAttr mlir::spirv::getStrArrayAttrForEnumList	(	Builder &	builder,
		ArrayRef< Ty >	enumValues,
		function_ref< StringRef(Ty)>	stringifyFn
	)

Definition at line 32 of file SPIRVParsingUtils.h.

References mlir::Builder::getStrArrayAttr().

◆ getTargetEnvAttrName()

StringRef mlir::spirv::getTargetEnvAttrName ( )

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

Definition at line 169 of file TargetAndABI.cpp.

Referenced by lookupTargetEnv().

◆ getUnaryOpResultType()

static Type mlir::spirv::getUnaryOpResultType ( Type operandType )

static

Result of a logical op must be a scalar or vector of boolean type.

Definition at line 41 of file SPIRVOpDefinition.cpp.

References mlir::get(), mlir::Type::getContext(), and mlir::Builder::getIntegerType().

◆ getVulkanElementPtr()

Value mlir::spirv::getVulkanElementPtr	(	const SPIRVTypeConverter &	typeConverter,
		MemRefType	baseType,
		Value	basePtr,
		ValueRange	indices,
		Location	loc,
		OpBuilder &	builder
	)

Definition at line 1221 of file SPIRVConversion.cpp.

References mlir::OpBuilder::create(), mlir::SPIRVTypeConverter::getIndexType(), getZero(), and mlir::affine::linearizeIndex().

Referenced by getElementPtr().

◆ hash_value()

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

Definition at line 1067 of file SPIRVTypes.cpp.

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

◆ hasOneBranchOpTo()

static bool mlir::spirv::hasOneBranchOpTo	(	Block &	srcBlock,
		Block &	dstBlock
	)

static

Returns true if the given srcBlock contains only one spirv.Branch to the given dstBlock.

Definition at line 247 of file ControlFlowOps.cpp.

References mlir::Block::back().

◆ hasOtherMerge()

static bool mlir::spirv::hasOtherMerge ( Region & region )

static

Returns true if a spirv.mlir.merge op outside the merge block.

Definition at line 262 of file ControlFlowOps.cpp.

References mlir::Region::empty(), and mlir::Region::getOps().

◆ isDirectInModuleLikeOp()

static bool mlir::spirv::isDirectInModuleLikeOp ( Operation * op )

static

Returns true if the given op is an module-like op that maintains a symbol table.

Definition at line 36 of file SPIRVOpDefinition.cpp.

References mlir::Operation::hasTrait().

◆ isMergeBlock()

static bool mlir::spirv::isMergeBlock ( Block & block )

static

Returns true if the given block only contains one spirv.mlir.merge op.

Definition at line 257 of file ControlFlowOps.cpp.

References mlir::Block::front().

◆ isNestedInFunctionOpInterface()

static bool mlir::spirv::isNestedInFunctionOpInterface ( Operation * op )

static

Returns true if the given op is a function-like op or nested in a function-like op without a module-like op in the middle.

Definition at line 24 of file SPIRVOpDefinition.cpp.

References mlir::Operation::getParentOp(), and mlir::Operation::hasTrait().

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

References mlir::OpBuilder::createOrFold(), mlir::detail::enumerate(), and mlir::get().

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

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

Referenced by lookupLocalWorkGroupSize().

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

References lookupEntryPointABI().

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

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

Referenced by getInterfaceVariables(), lookupTargetEnvOrDefault(), and lowerEntryPointABIAttr().

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

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

◆ mapMemorySpaceToOpenCLStorageClass()

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 119 of file MapMemRefStorageClassPass.cpp.

References OPENCL_STORAGE_SPACE_MAP_LIST, and STORAGE_SPACE_MAP_FN.

◆ mapMemorySpaceToVulkanStorageClass()

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 65 of file MapMemRefStorageClassPass.cpp.

References STORAGE_SPACE_MAP_FN, and VULKAN_STORAGE_SPACE_MAP_LIST.

◆ mapOpenCLStorageClassToMemorySpace()

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 146 of file MapMemRefStorageClassPass.cpp.

References OPENCL_STORAGE_SPACE_MAP_LIST, and STORAGE_SPACE_MAP_FN.

◆ mapVulkanStorageClassToMemorySpace()

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 92 of file MapMemRefStorageClassPass.cpp.

References STORAGE_SPACE_MAP_FN, and VULKAN_STORAGE_SPACE_MAP_LIST.

◆ needsInterfaceVarABIAttrs()

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

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

Referenced by getDefaultABIAttrs().

◆ parseControlAttribute()

template<typename EnumAttrClass , typename EnumClass >

static ParseResult mlir::spirv::parseControlAttribute	(	OpAsmParser &	parser,
		OperationState &	state,
		StringRef	attrName = `spirv::attributeName<EnumClass>()`
	)

static

Parses Function, Selection and Loop control attributes.

If no control is specified, "None" is used as a default.

Definition at line 29 of file ControlFlowOps.cpp.

◆ parseEnumKeywordAttr() [1/2]

template<typename EnumAttrClass , typename EnumClass = typename EnumAttrClass::ValueType>

ParseResult mlir::spirv::parseEnumKeywordAttr	(	EnumClass &	value,
		OpAsmParser &	parser,
		OperationState &	state,
		StringRef	attrName = `spirv::attributeName<EnumClass>()`
	)

Parses the next keyword in parser as an enumerant of the given EnumClass and inserts the enumerant into state as an 32-bit integer attribute with the enum class's name as attribute name.

Definition at line 113 of file SPIRVParsingUtils.h.

◆ parseEnumKeywordAttr() [2/2]

template<typename EnumClass , typename ParserType >

ParseResult mlir::spirv::parseEnumKeywordAttr	(	EnumClass &	value,
		ParserType &	parser,
		StringRef	attrName = `spirv::attributeName<EnumClass>()`
	)

Parses the next keyword in parser as an enumerant of the given EnumClass.

Definition at line 49 of file SPIRVParsingUtils.h.

Referenced by parseCooperativeMatrixType().

◆ parseEnumStrAttr() [1/2]

template<typename EnumAttrClass , typename EnumClass = typename EnumAttrClass::ValueType>

ParseResult mlir::spirv::parseEnumStrAttr	(	EnumClass &	value,
		OpAsmParser &	parser,
		OperationState &	state,
		StringRef	attrName = `spirv::attributeName<EnumClass>()`
	)

Parses the next string attribute in parser as an enumerant of the given EnumClass and inserts the enumerant into state as an 32-bit integer attribute with the enum class's name as attribute name.

Definition at line 97 of file SPIRVParsingUtils.h.

◆ parseEnumStrAttr() [2/2]

template<typename EnumClass >

ParseResult mlir::spirv::parseEnumStrAttr	(	EnumClass &	value,
		OpAsmParser &	parser,
		StringRef	attrName = `spirv::attributeName<EnumClass>()`
	)

Parses the next string attribute in parser as an enumerant of the given EnumClass.

Definition at line 70 of file SPIRVParsingUtils.h.

References mlir::AsmParser::emitError(), mlir::AsmParser::getBuilder(), mlir::AsmParser::getCurrentLocation(), mlir::Builder::getNoneType(), and mlir::AsmParser::parseAttribute().

Referenced by parseImageOperands().

◆ parseImageOperands()

static ParseResult mlir::spirv::parseImageOperands	(	OpAsmParser &	parser,
		spirv::ImageOperandsAttr &	attr
	)

static

Definition at line 49 of file SPIRVOpDefinition.cpp.

References mlir::get(), mlir::AsmParser::getContext(), parseEnumStrAttr(), mlir::AsmParser::parseOptionalLSquare(), and mlir::AsmParser::parseRSquare().

◆ parseMemoryAccessAttributes()

template<typename MemoryOpTy >

ParseResult mlir::spirv::parseMemoryAccessAttributes	(	OpAsmParser &	parser,
		OperationState &	state
	)

Parses optional memory access (a.k.a.

memory operand) attributes attached to a memory access operand/pointer. Specifically, parses the following syntax: ([ memory-access ])? where: memory-access ::= "None" | "Volatile" | "Aligned", integer-literal | "NonTemporal"

Definition at line 35 of file MemoryOps.cpp.

◆ parseSourceMemoryAccessAttributes()

template<typename MemoryOpTy >

static ParseResult mlir::spirv::parseSourceMemoryAccessAttributes	(	OpAsmParser &	parser,
		OperationState &	state
	)

static

Definition at line 70 of file MemoryOps.cpp.

◆ parseVariableDecorations()

ParseResult mlir::spirv::parseVariableDecorations	(	OpAsmParser &	parser,
		OperationState &	state
	)

Definition at line 21 of file SPIRVParsingUtils.cpp.

◆ populateSPIRVExpandExtendedMultiplicationPatterns()

void mlir::spirv::populateSPIRVExpandExtendedMultiplicationPatterns ( RewritePatternSet & patterns )

Appends patterns to expand extended multiplication and adition ops into regular arithmetic ops.

Extended arithmetic ops are not supported by the WebGPU Shading Language (WGSL).

Definition at line 248 of file SPIRVWebGPUTransforms.cpp.

References mlir::patterns.

◆ populateSPIRVExpandNonFiniteArithmeticPatterns()

void mlir::spirv::populateSPIRVExpandNonFiniteArithmeticPatterns ( RewritePatternSet & patterns )

Appends patterns to expand non-finite arithmetic ops IsNan and IsInf.

These are not supported by the WebGPU Shading Language (WGSL). We follow fast math assumptions and assume that all floating point values are finite.

Definition at line 256 of file SPIRVWebGPUTransforms.cpp.

References mlir::patterns.

◆ populateSPIRVGLCanonicalizationPatterns()

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

◆ printAccessChain()

template<typename Op >

static void mlir::spirv::printAccessChain	(	Op	op,
		ValueRange	indices,
		OpAsmPrinter &	printer
	)

static

Definition at line 324 of file MemoryOps.cpp.

References mlir::ValueRange::getTypes().

◆ printImageOperands()

static void mlir::spirv::printImageOperands	(	OpAsmPrinter &	printer,
		Operation *	imageOp,
		spirv::ImageOperandsAttr	attr
	)

static

Definition at line 64 of file SPIRVOpDefinition.cpp.

◆ printMemoryAccessAttribute()

template<typename MemoryOpTy >

static void mlir::spirv::printMemoryAccessAttribute	(	MemoryOpTy	memoryOp,
		OpAsmPrinter &	printer,
		SmallVectorImpl< StringRef > &	elidedAttrs,
		std::optional< spirv::MemoryAccess >	memoryAccessAtrrValue = `std::nullopt`,
		std::optional< uint32_t >	alignmentAttrValue = `std::nullopt`
	)

static

Definition at line 135 of file MemoryOps.cpp.

◆ printSourceMemoryAccessAttribute()

template<typename MemoryOpTy >

static void mlir::spirv::printSourceMemoryAccessAttribute	(	MemoryOpTy	memoryOp,
		OpAsmPrinter &	printer,
		SmallVectorImpl< StringRef > &	elidedAttrs,
		std::optional< spirv::MemoryAccess >	memoryAccessAtrrValue = `std::nullopt`,
		std::optional< uint32_t >	alignmentAttrValue = `std::nullopt`
	)

static

Definition at line 106 of file MemoryOps.cpp.

◆ printVariableDecorations()

void mlir::spirv::printVariableDecorations	(	Operation *	op,
		OpAsmPrinter &	printer,
		SmallVectorImpl< StringRef > &	elidedAttrs
	)

Definition at line 94 of file SPIRVOps.cpp.

References bindingName(), descriptorSetName(), mlir::Operation::getAttrOfType(), mlir::Operation::getAttrs(), and mlir::OpAsmPrinter::printOptionalAttrDict().

◆ processDecorationList()

template<typename AttrTy , typename EmitF >

LogicalResult mlir::spirv::processDecorationList	(	Location	loc,
		Decoration	decoration,
		Attribute	attrList,
		StringRef	attrName,
		EmitF	emitter
	)

Definition at line 230 of file Serializer.cpp.

References mlir::emitError().

◆ registerSPIRVTargetInterfaceExternalModels() [1/2]

void mlir::spirv::registerSPIRVTargetInterfaceExternalModels ( DialectRegistry & registry )

Registers the TargetAttrInterface for the #spirv.target_env attribute in the given registry.

Definition at line 44 of file Target.cpp.

References mlir::DialectRegistry::addExtension().

Referenced by mlir::registerAllDialects(), and registerSPIRVTargetInterfaceExternalModels().

◆ registerSPIRVTargetInterfaceExternalModels() [2/2]

void mlir::spirv::registerSPIRVTargetInterfaceExternalModels ( MLIRContext & context )

Registers the TargetAttrInterface for the #spirv.target_env attribute in the registry associated with the given context.

Definition at line 51 of file Target.cpp.

References mlir::MLIRContext::appendDialectRegistry(), and registerSPIRVTargetInterfaceExternalModels().

◆ 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::CopyMemoryOp >()

template<>

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

Definition at line 715 of file SerializeOps.cpp.

References encodeInstructionInto().

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

template<>

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

Definition at line 624 of file SerializeOps.cpp.

References encodeInstructionInto(), and encodeStringLiteralInto().

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

template<>

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

Definition at line 659 of file SerializeOps.cpp.

References encodeInstructionInto().

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

template<>

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

Definition at line 688 of file SerializeOps.cpp.

References encodeInstructionInto().

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

template<>

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

Definition at line 764 of file SerializeOps.cpp.

References encodeInstructionInto().

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

template<>

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

Definition at line 384 of file Serializer.cpp.

References mlir::spirv::ArrayType::getArrayStride().

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

template<>

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

Definition at line 394 of file Serializer.cpp.

References mlir::spirv::RuntimeArrayType::getArrayStride().

◆ stringifyTypeName()

template<typename T >

static StringRef mlir::spirv::stringifyTypeName ( )

static

◆ stringifyTypeName< FloatType >()

template<>

StringRef mlir::spirv::stringifyTypeName< FloatType > ( )

Definition at line 31 of file AtomicOps.cpp.

◆ stringifyTypeName< IntegerType >()

template<>

StringRef mlir::spirv::stringifyTypeName< IntegerType > ( )

Definition at line 26 of file AtomicOps.cpp.

◆ unrollVectorsInFuncBodies()

LogicalResult mlir::spirv::unrollVectorsInFuncBodies ( Operation * op )

Definition at line 1361 of file SPIRVConversion.cpp.

References mlir::applyPatternsGreedily(), mlir::Operation::getContext(), getNativeVectorShape(), options, mlir::patterns, mlir::vector::populateCastAwayVectorLeadingOneDimPatterns(), mlir::vector::populateVectorInsertExtractStridedSliceDecompositionPatterns(), mlir::vector::populateVectorShapeCastLoweringPatterns(), mlir::vector::populateVectorTransposeLoweringPatterns(), mlir::vector::populateVectorUnrollPatterns(), and mlir::vector::UnrollVectorOptions::setNativeShapeFn().

◆ unrollVectorsInSignatures()

LogicalResult mlir::spirv::unrollVectorsInSignatures ( Operation * op )

Definition at line 1348 of file SPIRVConversion.cpp.

References mlir::applyPatternsGreedily(), mlir::config, mlir::ExistingOps, mlir::Operation::getContext(), mlir::patterns, mlir::populateFuncOpVectorRewritePatterns(), and mlir::populateReturnOpVectorRewritePatterns().

◆ verifyAccessChain()

template<typename Op >

static LogicalResult mlir::spirv::verifyAccessChain	(	Op	accessChainOp,
		ValueRange	indices
	)

static

Definition at line 330 of file MemoryOps.cpp.

References mlir::OpState::emitOpError(), getElementPtrType(), and mlir::OpState::getLoc().

◆ verifyAtomicUpdateOp()

template<typename AtomicOpTy , typename ExpectedElementType >

static LogicalResult mlir::spirv::verifyAtomicUpdateOp ( Operation * op )

static

Definition at line 37 of file AtomicOps.cpp.

References mlir::Operation::emitOpError(), mlir::Operation::getAttrOfType(), mlir::Operation::getName(), mlir::Operation::getOperand(), mlir::Value::getType(), and verifyMemorySemantics().

◆ verifyCastOp()

static LogicalResult mlir::spirv::verifyCastOp	(	Operation *	op,
		bool	requireSameBitWidth = `true`,
		bool	skipBitWidthCheck = `false`
	)

static

Definition at line 24 of file CastOps.cpp.

References mlir::Operation::emitOpError(), mlir::Operation::getOperand(), mlir::Operation::getResult(), and mlir::Value::getType().

◆ verifyCoopMatrixAccess()

static LogicalResult mlir::spirv::verifyCoopMatrixAccess	(	Operation *	op,
		Type	pointer,
		Type	coopMatrix,
		spirv::MemoryAccessAttr	memoryOperand
	)

static

Definition at line 25 of file CooperativeMatrixOps.cpp.

References mlir::Operation::emitOpError().

◆ verifyGroupNonUniformArithmeticOp()

template<typename OpTy >

static LogicalResult mlir::spirv::verifyGroupNonUniformArithmeticOp ( Operation * groupOp )

static

Definition at line 24 of file GroupOps.cpp.

References mlir::Operation::emitOpError(), extractValueFromConstOp(), mlir::Operation::getAttrOfType(), mlir::Value::getDefiningOp(), mlir::Operation::getName(), mlir::Operation::getNumOperands(), and mlir::Operation::getOperand().

◆ verifyGroupNonUniformShuffleOp()

template<typename OpTy >

static LogicalResult mlir::spirv::verifyGroupNonUniformShuffleOp ( OpTy op )

static

Definition at line 143 of file GroupOps.cpp.

◆ verifyGroupOp()

template<typename Op >

static LogicalResult mlir::spirv::verifyGroupOp ( Op op )

static

Definition at line 335 of file GroupOps.cpp.

References mlir::OpState::emitOpError().

◆ verifyIntegerDotProduct()

template<typename IntegerDotProductOpTy >

static LogicalResult mlir::spirv::verifyIntegerDotProduct ( Operation * op )

static

Definition at line 29 of file IntegerDotProductOps.cpp.

References mlir::Operation::emitOpError(), mlir::ValueTypeRange< ValueRangeT >::front(), mlir::Operation::getAttr(), getBitWidth(), mlir::Operation::getName(), mlir::Operation::getNumOperands(), mlir::Operation::getNumResults(), mlir::Operation::getOperand(), mlir::Operation::getResultTypes(), mlir::Value::getType(), and mlir::Operation::hasAttr().

◆ verifyLoadStorePtrAndValTypes()

template<typename LoadStoreOpTy >

static LogicalResult mlir::spirv::verifyLoadStorePtrAndValTypes	(	LoadStoreOpTy	op,
		Value	ptr,
		Value	val
	)

static

Definition at line 161 of file MemoryOps.cpp.

References mlir::Value::getType().

◆ verifyMemoryAccessAttribute()

template<typename MemoryOpTy >

static LogicalResult mlir::spirv::verifyMemoryAccessAttribute ( MemoryOpTy memoryOp )

static

Definition at line 176 of file MemoryOps.cpp.

◆ verifyMemorySemantics()

LogicalResult mlir::spirv::verifyMemorySemantics	(	Operation *	op,
		spirv::MemorySemantics	memorySemantics
	)

Definition at line 70 of file SPIRVOps.cpp.

References mlir::Operation::emitError().

Referenced by verifyAtomicUpdateOp().

◆ verifySourceMemoryAccessAttribute()

template<typename MemoryOpTy >

static LogicalResult mlir::spirv::verifySourceMemoryAccessAttribute ( MemoryOpTy memoryOp )

static

Definition at line 220 of file MemoryOps.cpp.

Variable Documentation

◆ kGeneratorNumber

constexpr uint32_t mlir::spirv::kGeneratorNumber = 22

constexpr

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

constexpr

SPIR-V binary header word count.

Definition at line 25 of file SPIRVBinaryUtils.h.

Referenced by mlir::spirv::Serializer::collect().

◆ kMagicNumber

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

constexpr

SPIR-V magic number.

Definition at line 28 of file SPIRVBinaryUtils.h.

Referenced by appendModuleHeader().

Namespaces

Classes

Typedefs

Functions

Variables

Typedef Documentation

◆ BlockMergeInfoMap

◆ GetTargetEnvFn

◆ MemorySpaceToStorageClassMap

◆ SymbolRenameListener

Function Documentation

◆ appendModuleHeader()

◆ combine()

◆ convertMemRefTypesAndAttrs()

◆ createUnifyAliasedResourcePass()

◆ decodeStringLiteral()

◆ deserialize()

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

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

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

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

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

◆ encodeInstructionInto()

◆ encodeStringLiteralInto()

◆ extractValueFromConstOp()

◆ getAddressingModel()

◆ getBitWidth()

◆ getBuiltinVariableValue()

◆ getComputeVectorSize()

◆ getDecorationName()

◆ getDefaultResourceLimits()

◆ getDefaultTargetEnv()

◆ getDirectImpliedCapabilities()

◆ getElementPtr()

◆ getElementPtrType()

◆ getEntryPointABIAttr()

◆ getEntryPointABIAttrName()

◆ getExecutionModel()

◆ getImpliedExtensions()

◆ getIntegerDotProductCapabilities()

◆ getIntegerDotProductExtensions()

◆ getIntegerDotProductMaxVersion()

◆ getIntegerDotProductMinVersion()

◆ getInterfaceVarABIAttr()

◆ getInterfaceVarABIAttrName()

◆ getMemoryModel()

◆ getMemorySpaceToStorageClassTarget()

◆ getNativeVectorShape()

◆ getNativeVectorShapeImpl() [1/2]

◆ getNativeVectorShapeImpl() [2/2]

◆ getOpenCLElementPtr()

◆ getPrefixedOpcode()

◆ getPushConstantValue()

◆ getRecursiveImpliedCapabilities()

◆ getStrArrayAttrForEnumList()

◆ getTargetEnvAttrName()

◆ getUnaryOpResultType()

◆ getVulkanElementPtr()

◆ hash_value()

◆ hasOneBranchOpTo()

◆ hasOtherMerge()

◆ isDirectInModuleLikeOp()

◆ isMergeBlock()

◆ isNestedInFunctionOpInterface()

◆ linearizeIndex()

◆ lookupEntryPointABI()

◆ lookupLocalWorkGroupSize()

◆ lookupTargetEnv()

◆ lookupTargetEnvOrDefault()

◆ mapMemorySpaceToOpenCLStorageClass()

◆ mapMemorySpaceToVulkanStorageClass()

◆ mapOpenCLStorageClassToMemorySpace()

◆ mapVulkanStorageClassToMemorySpace()

◆ needsInterfaceVarABIAttrs()

◆ parseControlAttribute()

◆ parseEnumKeywordAttr() [1/2]

◆ parseEnumKeywordAttr() [2/2]

◆ parseEnumStrAttr() [1/2]

◆ parseEnumStrAttr() [2/2]

◆ parseImageOperands()