MLIR
18.0.0git
|
#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"
Go to the source code of this file.
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" |
#define CLIENT_MAP | ( | client, | |
storage | |||
) |
#define DEBUG_TYPE "spirv-to-llvm-pattern" |
Definition at line 28 of file SPIRVToLLVM.cpp.
#define DISPATCH | ( | functionControl, | |
llvmAttr | |||
) |
#define STORAGE_SPACE_MAP | ( | storage, | |
space | |||
) |
|
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().
|
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().
|
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().
|
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().
|
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().
|
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().
|
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().
|
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().
|
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().
|
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().
|
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().
|
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.
|
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().
|
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().
|
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().
|
static |
Definition at line 302 of file SPIRVToLLVM.cpp.
References CLIENT_MAP, defaultAddressSpace, and mlir::gpu::amd::OpenCL.
Referenced by convertPointerType().
|
static |
Definition at line 280 of file SPIRVToLLVM.cpp.
References defaultAddressSpace, and STORAGE_SPACE_MAP.
Creates IntegerAttribute
with all bits set for given type.
Definition at line 93 of file SPIRVToLLVM.cpp.
References mlir::Builder::getIntegerAttr().
Referenced by createConstantAllBitsSet().
|
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().
|
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().
|
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:
Base
Offset
or Count
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().
|
static |
Utility for spirv.Load
and spirv.Store
conversion.
Definition at line 238 of file SPIRVToLLVM.cpp.
|
constexpr |
Definition at line 36 of file SPIRVToLLVM.cpp.
Referenced by mapToAddressSpace(), and mapToOpenCLAddressSpace().
|
staticconstexpr |
Hook for descriptor set and binding number encoding.
Definition at line 1696 of file SPIRVToLLVM.cpp.
|
staticconstexpr |
Definition at line 1697 of file SPIRVToLLVM.cpp.