#include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVM.h"
#include "mlir/Conversion/LLVMCommon/Pattern.h"
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
#include "mlir/Dialect/SPIRV/Utils/LayoutUtils.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/Support/LogicalResult.h"
#include "mlir/Transforms/DialectConversion.h"
#include "llvm/Support/Debug.h"
#include "llvm/Support/FormatVariadic.h"

Macros
#define	DEBUG_TYPE "spirv-to-llvm-pattern"

#define	STORAGE_SPACE_MAP(storage, space)

#define	CLIENT_MAP(client, storage)

#define	DISPATCH(functionControl, llvmAttr)

Functions
static bool	isSignedIntegerOrVector (Type type)
	Returns true if the given type is a signed integer or vector type. More...

static bool	isUnsignedIntegerOrVector (Type type)
	Returns true if the given type is an unsigned integer or vector type. More...

static std::optional< uint64_t >	getIntegerOrVectorElementWidth (Type type)
	Returns the width of an integer or of the element type of an integer vector, if applicable. More...

static unsigned	getBitWidth (Type type)
	Returns the bit width of integer, float or vector of float or integer values. More...

static unsigned	getLLVMTypeBitWidth (Type type)
	Returns the bit width of LLVMType integer or vector. More...

static IntegerAttr	minusOneIntegerAttribute (Type type, Builder builder)
	Creates `IntegerAttribute` with all bits set for given type. More...

static Value	createConstantAllBitsSet (Location loc, Type srcType, Type dstType, PatternRewriter &rewriter)
	Creates `llvm.mlir.constant` with all bits set for the given type. More...

static Value	createFPConstant (Location loc, Type srcType, Type dstType, PatternRewriter &rewriter, double value)
	Creates `llvm.mlir.constant` with a floating-point scalar or vector value. More...

static Value	optionallyTruncateOrExtend (Location loc, Value value, Type llvmType, PatternRewriter &rewriter)
	Utility function for bitfield ops: More...

static Value	broadcast (Location loc, Value toBroadcast, unsigned numElements, LLVMTypeConverter &typeConverter, ConversionPatternRewriter &rewriter)
	Broadcasts the value to vector with `numElements` number of elements. More...

static Value	optionallyBroadcast (Location loc, Value value, Type srcType, LLVMTypeConverter &typeConverter, ConversionPatternRewriter &rewriter)
	Broadcasts the value. If `srcType` is a scalar, the value remains unchanged. More...

static Value	processCountOrOffset (Location loc, Value value, Type srcType, Type dstType, LLVMTypeConverter &converter, ConversionPatternRewriter &rewriter)
	Utility function for bitfield ops: `BitFieldInsert`, `BitFieldSExtract` and `BitFieldUExtract`. More...

static std::optional< Type >	convertStructTypeWithOffset (spirv::StructType type, LLVMTypeConverter &converter)
	Converts SPIR-V struct with a regular (according to `VulkanLayoutUtils`) offset to LLVM struct. More...

static Type	convertStructTypePacked (spirv::StructType type, LLVMTypeConverter &converter)
	Converts SPIR-V struct with no offset to packed LLVM struct. More...

static Value	createI32ConstantOf (Location loc, PatternRewriter &rewriter, unsigned value)
	Creates LLVM dialect constant with the given value. More...

static LogicalResult	replaceWithLoadOrStore (Operation *op, ValueRange operands, ConversionPatternRewriter &rewriter, LLVMTypeConverter &typeConverter, unsigned alignment, bool isVolatile, bool isNonTemporal)
	Utility for `spirv.Load` and `spirv.Store` conversion. More...

static std::optional< Type >	convertArrayType (spirv::ArrayType type, TypeConverter &converter)
	Converts SPIR-V array type to LLVM array. More...

static unsigned	mapToOpenCLAddressSpace (spirv::StorageClass storageClass)

static unsigned	mapToAddressSpace (spirv::ClientAPI clientAPI, spirv::StorageClass storageClass)

static Type	convertPointerType (spirv::PointerType type, LLVMTypeConverter &converter, spirv::ClientAPI clientAPI)
	Converts SPIR-V pointer type to LLVM pointer. More...

static std::optional< Type >	convertRuntimeArrayType (spirv::RuntimeArrayType type, TypeConverter &converter)
	Converts SPIR-V runtime array to LLVM array. More...

static std::optional< Type >	convertStructType (spirv::StructType type, LLVMTypeConverter &converter)
	Converts SPIR-V struct to LLVM struct. More...

Variables
constexpr unsigned	defaultAddressSpace = 0

static constexpr StringRef	kBinding = "binding"
	Hook for descriptor set and binding number encoding. More...

static constexpr StringRef	kDescriptorSet = "descriptor_set"

Macro Definition Documentation

◆ CLIENT_MAP

#define CLIENT_MAP	(	client,
		storage
	)

Value:

case spirv::ClientAPI::client: \

return mapTo##client##AddressSpace(storage);

◆ DEBUG_TYPE

#define DEBUG_TYPE "spirv-to-llvm-pattern"

Definition at line 28 of file SPIRVToLLVM.cpp.

◆ DISPATCH

#define DISPATCH	(	functionControl,
		llvmAttr
	)

Value:

  case functionControl:                                                        \
    newFuncOp->setAttr("passthrough", ArrayAttr::get(context, {llvmAttr}));    \
    break;

◆ STORAGE_SPACE_MAP

#define STORAGE_SPACE_MAP	(	storage,
		space
	)

Value:

case spirv::StorageClass::storage: \

return space;

Function Documentation

◆ broadcast()

static Value broadcast	(	Location	loc,
		Value	toBroadcast,
		unsigned	numElements,
		LLVMTypeConverter &	typeConverter,
		ConversionPatternRewriter &	rewriter
	)

static

Broadcasts the value to vector with numElements number of elements.

Definition at line 157 of file SPIRVToLLVM.cpp.

References mlir::LLVMTypeConverter::convertType(), mlir::OpBuilder::create(), mlir::get(), mlir::Builder::getI32IntegerAttr(), mlir::Builder::getIntegerType(), and mlir::Value::getType().

Referenced by optionallyBroadcast().

◆ convertArrayType()

static std::optional<Type> convertArrayType	(	spirv::ArrayType	type,
		TypeConverter &	converter
	)

static

Converts SPIR-V array type to LLVM array.

Natural stride (according to VulkanLayoutUtils) is also mapped to LLVM array. This has to be respected when converting ops that manipulate array types.

Definition at line 267 of file SPIRVToLLVM.cpp.

References mlir::TypeConverter::convertType(), mlir::get(), mlir::spirv::ArrayType::getArrayStride(), mlir::spirv::ArrayType::getElementType(), and mlir::spirv::ArrayType::getNumElements().

Referenced by mlir::populateSPIRVToLLVMTypeConversion().

◆ convertPointerType()

static Type convertPointerType	(	spirv::PointerType	type,
		LLVMTypeConverter &	converter,
		spirv::ClientAPI	clientAPI
	)

static

Converts SPIR-V pointer type to LLVM pointer.

Pointer's storage class is not modelled at the moment.

Definition at line 317 of file SPIRVToLLVM.cpp.

References mlir::LLVMTypeConverter::convertType(), mlir::spirv::PointerType::getPointeeType(), mlir::LLVMTypeConverter::getPointerType(), mlir::spirv::PointerType::getStorageClass(), and mapToAddressSpace().

Referenced by mlir::populateSPIRVToLLVMTypeConversion().

◆ convertRuntimeArrayType()

static std::optional<Type> convertRuntimeArrayType	(	spirv::RuntimeArrayType	type,
		TypeConverter &	converter
	)

static

Converts SPIR-V runtime array to LLVM array.

Since LLVM allows indexing over the bounds, the runtime array is converted to a 0-sized LLVM array. There is no modelling of array stride at the moment.

Definition at line 328 of file SPIRVToLLVM.cpp.

References mlir::TypeConverter::convertType(), mlir::get(), mlir::spirv::RuntimeArrayType::getArrayStride(), and mlir::spirv::RuntimeArrayType::getElementType().

Referenced by mlir::populateSPIRVToLLVMTypeConversion().

◆ convertStructType()

static std::optional<Type> convertStructType	(	spirv::StructType	type,
		LLVMTypeConverter &	converter
	)

static

Converts SPIR-V struct to LLVM struct.

There is no support of structs with member decorations. Also, only natural offset is supported.

Definition at line 338 of file SPIRVToLLVM.cpp.

References convertStructTypePacked(), convertStructTypeWithOffset(), mlir::spirv::StructType::getMemberDecorations(), and mlir::spirv::StructType::hasOffset().

Referenced by mlir::populateSPIRVToLLVMTypeConversion().

◆ convertStructTypePacked()

static Type convertStructTypePacked	(	spirv::StructType	type,
		LLVMTypeConverter &	converter
	)

static

Converts SPIR-V struct with no offset to packed LLVM struct.

Definition at line 219 of file SPIRVToLLVM.cpp.

References mlir::spirv::StructType::getElementTypes(), and mlir::LLVM::LLVMStructType::getLiteral().

Referenced by convertStructType().

◆ convertStructTypeWithOffset()

static std::optional<Type> convertStructTypeWithOffset	(	spirv::StructType	type,
		LLVMTypeConverter &	converter
	)

static

Converts SPIR-V struct with a regular (according to VulkanLayoutUtils) offset to LLVM struct.

Otherwise, the conversion is not supported.

Definition at line 205 of file SPIRVToLLVM.cpp.

References mlir::VulkanLayoutUtils::decorateType(), mlir::spirv::StructType::getElementTypes(), and mlir::LLVM::LLVMStructType::getLiteral().

Referenced by convertStructType().

◆ createConstantAllBitsSet()

static Value createConstantAllBitsSet	(	Location	loc,
		Type	srcType,
		Type	dstType,
		PatternRewriter &	rewriter
	)

static

Creates llvm.mlir.constant with all bits set for the given type.

Definition at line 103 of file SPIRVToLLVM.cpp.

References mlir::OpBuilder::create(), mlir::DenseElementsAttr::get(), and minusOneIntegerAttribute().

◆ createFPConstant()

static Value createFPConstant	(	Location	loc,
		Type	srcType,
		Type	dstType,
		PatternRewriter &	rewriter,
		double	value
	)

static

Creates llvm.mlir.constant with a floating-point scalar or vector value.

Definition at line 116 of file SPIRVToLLVM.cpp.

References mlir::OpBuilder::create(), mlir::DenseElementsAttr::get(), and mlir::Builder::getFloatAttr().

◆ createI32ConstantOf()

static Value createI32ConstantOf	(	Location	loc,
		PatternRewriter &	rewriter,
		unsigned	value
	)

static

Creates LLVM dialect constant with the given value.

Definition at line 230 of file SPIRVToLLVM.cpp.

References mlir::OpBuilder::create(), mlir::get(), mlir::Builder::getContext(), mlir::Builder::getI32Type(), and mlir::Builder::getIntegerAttr().

◆ getBitWidth()

static unsigned getBitWidth ( Type type )

static

Returns the bit width of integer, float or vector of float or integer values.

Definition at line 72 of file SPIRVToLLVM.cpp.

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

Referenced by mlir::ConstantIntRanges::getConstantValue(), mlir::ConstantIntRanges::intersection(), optionallyTruncateOrExtend(), and mlir::ConstantIntRanges::rangeUnion().

◆ getIntegerOrVectorElementWidth()

static std::optional<uint64_t> getIntegerOrVectorElementWidth ( Type type )

static

Returns the width of an integer or of the element type of an integer vector, if applicable.

Definition at line 62 of file SPIRVToLLVM.cpp.

◆ getLLVMTypeBitWidth()

static unsigned getLLVMTypeBitWidth ( Type type )

static

Returns the bit width of LLVMType integer or vector.

Definition at line 85 of file SPIRVToLLVM.cpp.

References mlir::LLVM::getVectorElementType(), and mlir::LLVM::isCompatibleVectorType().

Referenced by optionallyTruncateOrExtend().

◆ isSignedIntegerOrVector()

static bool isSignedIntegerOrVector ( Type type )

static

Returns true if the given type is a signed integer or vector type.

Definition at line 43 of file SPIRVToLLVM.cpp.

References mlir::Type::isSignedInteger().

◆ isUnsignedIntegerOrVector()

static bool isUnsignedIntegerOrVector ( Type type )

static

Returns true if the given type is an unsigned integer or vector type.

Definition at line 52 of file SPIRVToLLVM.cpp.

References mlir::Type::isUnsignedInteger().

◆ mapToAddressSpace()

static unsigned mapToAddressSpace	(	spirv::ClientAPI	clientAPI,
		spirv::StorageClass	storageClass
	)

static

Definition at line 302 of file SPIRVToLLVM.cpp.

References CLIENT_MAP, defaultAddressSpace, and mlir::gpu::amd::OpenCL.

Referenced by convertPointerType().

◆ mapToOpenCLAddressSpace()

static unsigned mapToOpenCLAddressSpace ( spirv::StorageClass storageClass )

static

Definition at line 280 of file SPIRVToLLVM.cpp.

References defaultAddressSpace, and STORAGE_SPACE_MAP.

◆ minusOneIntegerAttribute()

static IntegerAttr minusOneIntegerAttribute	(	Type	type,
		Builder	builder
	)

static

Creates IntegerAttribute with all bits set for given type.

Definition at line 93 of file SPIRVToLLVM.cpp.

References mlir::Builder::getIntegerAttr().

Referenced by createConstantAllBitsSet().

◆ optionallyBroadcast()

static Value optionallyBroadcast	(	Location	loc,
		Value	value,
		Type	srcType,
		LLVMTypeConverter &	typeConverter,
		ConversionPatternRewriter &	rewriter
	)

static

Broadcasts the value. If srcType is a scalar, the value remains unchanged.

Definition at line 174 of file SPIRVToLLVM.cpp.

References broadcast().

Referenced by processCountOrOffset().

◆ optionallyTruncateOrExtend()

static Value optionallyTruncateOrExtend	(	Location	loc,
		Value	value,
		Type	llvmType,
		PatternRewriter &	rewriter
	)

static

Utility function for bitfield ops:

BitFieldInsert
BitFieldSExtract
BitFieldUExtract Truncates or extends the value. If the bitwidth of the value is the same as llvmType bitwidth, the value remains unchanged.

Definition at line 136 of file SPIRVToLLVM.cpp.

References mlir::OpBuilder::create(), getBitWidth(), getLLVMTypeBitWidth(), mlir::Value::getType(), and mlir::LLVM::isCompatibleType().

Referenced by processCountOrOffset().

◆ processCountOrOffset()

static Value processCountOrOffset	(	Location	loc,
		Value	value,
		Type	srcType,
		Type	dstType,
		LLVMTypeConverter &	converter,
		ConversionPatternRewriter &	rewriter
	)

static

Utility function for bitfield ops: BitFieldInsert, BitFieldSExtract and BitFieldUExtract.

Broadcast Offset and Count to match the type of Base. If Base is of a vector type, construct a vector that has:

same number of elements as Base
each element has the type that is the same as the type of Offset or Count
each element has the same value as Offset or Count Then cast Offset and Count if their bit width is different from Base bit width.

Definition at line 194 of file SPIRVToLLVM.cpp.

References optionallyBroadcast(), and optionallyTruncateOrExtend().

◆ replaceWithLoadOrStore()

static LogicalResult replaceWithLoadOrStore	(	Operation *	op,
		ValueRange	operands,
		ConversionPatternRewriter &	rewriter,
		LLVMTypeConverter &	typeConverter,
		unsigned	alignment,
		bool	isVolatile,
		bool	isNonTemporal
	)

static

Utility for spirv.Load and spirv.Store conversion.

Definition at line 238 of file SPIRVToLLVM.cpp.

Variable Documentation

◆ defaultAddressSpace

constexpr unsigned defaultAddressSpace = 0

constexpr

Definition at line 36 of file SPIRVToLLVM.cpp.

Referenced by mapToAddressSpace(), and mapToOpenCLAddressSpace().

◆ kBinding

constexpr StringRef kBinding = "binding"

staticconstexpr

Hook for descriptor set and binding number encoding.

Definition at line 1696 of file SPIRVToLLVM.cpp.

◆ kDescriptorSet

constexpr StringRef kDescriptorSet = "descriptor_set"

staticconstexpr

Definition at line 1697 of file SPIRVToLLVM.cpp.

Macros

Functions

Variables

Macro Definition Documentation

◆ CLIENT_MAP

◆ DEBUG_TYPE

◆ DISPATCH

◆ STORAGE_SPACE_MAP

Function Documentation

◆ broadcast()

◆ convertArrayType()

◆ convertPointerType()

◆ convertRuntimeArrayType()

◆ convertStructType()

◆ convertStructTypePacked()

◆ convertStructTypeWithOffset()

◆ createConstantAllBitsSet()

◆ createFPConstant()

◆ createI32ConstantOf()

◆ getBitWidth()

◆ getIntegerOrVectorElementWidth()

◆ getLLVMTypeBitWidth()

◆ isSignedIntegerOrVector()

◆ isUnsignedIntegerOrVector()

◆ mapToAddressSpace()

◆ mapToOpenCLAddressSpace()

◆ minusOneIntegerAttribute()

◆ optionallyBroadcast()

◆ optionallyTruncateOrExtend()

◆ processCountOrOffset()

◆ replaceWithLoadOrStore()

Variable Documentation

◆ defaultAddressSpace

◆ kBinding

◆ kDescriptorSet