OSchip
/
llvm-project
mirror of https://github.com/THU-DSP-LAB/llvm-project.git
13
0
Fork 6
Code Issues Packages Projects Releases Wiki Activity

[mlir][spirv] Replace StructAttrs with AttrDefs 

Depends on D127370

Reviewed By: antiagainst

Differential Revision: https://reviews.llvm.org/D127373
This commit is contained in:
Mogball 2022-06-09 21:35:32 +00:00
parent b59c2315af
commit a31ff0af9b
48 changed files with 644 additions and 641 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
mlir/include/mlir/Dialect/SPIRV/IR/CMakeLists.txt
View File

@ -31,8 +31,8 @@ mlir_tablegen(SPIRVAttrUtils.inc -gen-spirv-attr-utils)
add_public_tablegen_target(MLIRSPIRVAttrUtilsGen)
add_dependencies(mlir-headers MLIRSPIRVAttrUtilsGen)
set(LLVM_TARGET_DEFINITIONS TargetAndABI.td)
mlir_tablegen(TargetAndABI.h.inc -gen-struct-attr-decls)
mlir_tablegen(TargetAndABI.cpp.inc -gen-struct-attr-defs)
add_public_tablegen_target(MLIRSPIRVTargetAndABIIncGen)
add_dependencies(mlir-headers MLIRSPIRVTargetAndABIIncGen)
set(LLVM_TARGET_DEFINITIONS SPIRVAttributes.td)
mlir_tablegen(SPIRVAttributes.h.inc -gen-attrdef-decls)
mlir_tablegen(SPIRVAttributes.cpp.inc -gen-attrdef-defs)
add_public_tablegen_target(MLIRSPIRVAttributeIncGen)
add_dependencies(mlir-headers MLIRSPIRVAttributeIncGen)

10
mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.h
View File

@ -18,7 +18,8 @@
#include "mlir/Support/LLVM.h"
// Pull in TableGen'erated SPIR-V attribute definitions for target and ABI.
#include "mlir/Dialect/SPIRV/IR/TargetAndABI.h.inc"
#define GET_ATTRDEF_CLASSES
#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h.inc"
namespace mlir {
namespace spirv {
@ -139,7 +140,7 @@ public:
/// Gets a TargetEnvAttr instance.
static TargetEnvAttr get(VerCapExtAttr triple, Vendor vendorID,
DeviceType deviceType, uint32_t deviceId,
DictionaryAttr limits);
ResourceLimitsAttr limits);
/// Returns the attribute kind's name (without the 'spv.' prefix).
static StringRef getKindName();
@ -171,11 +172,6 @@ public:
/// Returns the target resource limits.
ResourceLimitsAttr getResourceLimits() const;
static LogicalResult verify(function_ref<InFlightDiagnostic()> emitError,
VerCapExtAttr triple, Vendor vendorID,
DeviceType deviceType, uint32_t deviceID,
DictionaryAttr limits);
};
} // namespace spirv
} // namespace mlir

63
mlir/include/mlir/Dialect/SPIRV/IR/TargetAndABI.td → mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.td
View File

@ -23,12 +23,18 @@
include "mlir/Dialect/SPIRV/IR/SPIRVBase.td"
class SPV_Attr<string attrName, string attrMnemonic, list<Trait> traits = []>
: AttrDef<SPIRV_Dialect, attrName> {
let mnemonic = attrMnemonic;
}
// For entry functions, this attribute specifies information related to entry
// points in the generated SPIR-V module:
// 1) WorkGroup Size.
def SPV_EntryPointABIAttr : StructAttr<"EntryPointABIAttr", SPIRV_Dialect, [
StructFieldAttr<"local_size", OptionalAttr<I32ElementsAttr>>
]>;
def SPV_EntryPointABIAttr : SPV_Attr<"EntryPointABI", "entry_point_abi"> {
let parameters = (ins OptionalParameter<"DenseIntElementsAttr">:$local_size);
let assemblyFormat = "`<` struct(params) `>`";
}
def SPV_ExtensionArrayAttr : TypedArrayAttrBase<
SPV_ExtensionAttr, "SPIR-V extension array attribute">;
@ -40,16 +46,19 @@ def SPV_CapabilityArrayAttr : TypedArrayAttrBase<
// target. Represents `VkCooperativeMatrixPropertiesNV`. See
// https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkCooperativeMatrixPropertiesNV.html
def SPV_CooperativeMatrixPropertiesNVAttr :
StructAttr<"CooperativeMatrixPropertiesNVAttr", SPIRV_Dialect, [
StructFieldAttr<"m_size", I32Attr>,
StructFieldAttr<"n_size", I32Attr>,
StructFieldAttr<"k_size", I32Attr>,
StructFieldAttr<"a_type", TypeAttr>,
StructFieldAttr<"b_type", TypeAttr>,
StructFieldAttr<"c_type", TypeAttr>,
StructFieldAttr<"result_type", TypeAttr>,
StructFieldAttr<"scope", SPV_ScopeAttr>
]>;
SPV_Attr<"CooperativeMatrixPropertiesNV", "coop_matrix_props"> {
let parameters = (ins
"int":$m_size,
"int":$n_size,
"int":$k_size,
"mlir::Type":$a_type,
"mlir::Type":$b_type,
"mlir::Type":$c_type,
"mlir::Type":$result_type,
"mlir::spirv::ScopeAttr":$scope
);
let assemblyFormat = "`<` struct(params) `>`";
}
def SPV_CooperativeMatrixPropertiesNVArrayAttr :
TypedArrayAttrBase<SPV_CooperativeMatrixPropertiesNVAttr,
@ -63,28 +72,32 @@ def SPV_CooperativeMatrixPropertiesNVArrayAttr :
// The following ones are those affecting SPIR-V CodeGen. Their default value
// are the from Vulkan limit requirements:
// https://www.khronos.org/registry/vulkan/specs/1.2-extensions/html/vkspec.html#limits-minmax
def SPV_ResourceLimitsAttr : StructAttr<"ResourceLimitsAttr", SPIRV_Dialect, [
def SPV_ResourceLimitsAttr : SPV_Attr<"ResourceLimits", "resource_limits"> {
let parameters = (ins
// The maximum total storage size, in bytes, available for variables
// declared with the Workgroup storage class.
StructFieldAttr<"max_compute_shared_memory_size",
DefaultValuedAttr<I32Attr, "16384">>,
DefaultValuedParameter<"int", "16384">:$max_compute_shared_memory_size,
// The maximum total number of compute shader invocations in a single local
// workgroup.
StructFieldAttr<"max_compute_workgroup_invocations",
DefaultValuedAttr<I32Attr, "128">>,
DefaultValuedParameter<"int", "128">:$max_compute_workgroup_invocations,
// The maximum size of a local compute workgroup, per dimension.
StructFieldAttr<"max_compute_workgroup_size",
DefaultValuedAttr<I32ElementsAttr, "{128, 128, 64}">>,
DefaultValuedParameter<
"ArrayAttr",
"$_builder.getI32ArrayAttr({128, 128, 64})"
>:$max_compute_workgroup_size,
// The default number of invocations in each subgroup.
StructFieldAttr<"subgroup_size", DefaultValuedAttr<I32Attr, "32">>,
DefaultValuedParameter<"int", "32">:$subgroup_size,
// The configurations of cooperative matrix operations
// supported. Default is an empty list.
StructFieldAttr<
"cooperative_matrix_properties_nv",
DefaultValuedAttr<SPV_CooperativeMatrixPropertiesNVArrayAttr, "{}">>
]>;
DefaultValuedParameter<
"ArrayAttr",
"nullptr"
>:$cooperative_matrix_properties_nv
);
let assemblyFormat = "`<` struct(params) `>`";
}
#endif // MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI

8
mlir/include/mlir/Dialect/SPIRV/IR/SPIRVBase.td
View File

@ -48,7 +48,6 @@ def SPIRV_Dialect : Dialect {
let cppNamespace = "::mlir::spirv";
let useDefaultTypePrinterParser = 1;
let useDefaultAttributePrinterParser = 1;
let hasConstantMaterializer = 1;
let hasOperationAttrVerify = 1;
let hasRegionArgAttrVerify = 1;
@ -65,6 +64,13 @@ def SPIRV_Dialect : Dialect {
/// Returns the attribute name to use when specifying decorations on results
/// of operations.
static std::string getAttributeName(Decoration decoration);
/// Dialect attribute parsing hook.
Attribute parseAttribute(
DialectAsmParser &parser, Type type) const override;
/// Dialect attribute printing hook.
void printAttribute(
Attribute attr, DialectAsmPrinter &printer) const override;
}];
}

2
mlir/lib/Dialect/SPIRV/IR/CMakeLists.txt
View File

@ -16,13 +16,13 @@ add_mlir_dialect_library(MLIRSPIRV
${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/SPIRV
DEPENDS
MLIRSPIRVAttributeIncGen
MLIRSPIRVAttrUtilsGen
MLIRSPIRVAvailabilityIncGen
MLIRSPIRVCanonicalizationIncGen
MLIRSPIRVEnumAvailabilityIncGen
MLIRSPIRVEnumsIncGen
MLIRSPIRVOpsIncGen
MLIRSPIRVTargetAndABIIncGen
LINK_LIBS PUBLIC
MLIRControlFlowInterfaces

352
mlir/lib/Dialect/SPIRV/IR/SPIRVAttributes.cpp
View File

@ -10,8 +10,11 @@
#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVTypes.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/DialectImplementation.h"
#include "llvm/ADT/TypeSwitch.h"
using namespace mlir;
using namespace mlir::spirv;
//===----------------------------------------------------------------------===//
// TableGen'erated attribute utility functions
@ -21,15 +24,6 @@ namespace mlir {
namespace spirv {
#include "mlir/Dialect/SPIRV/IR/SPIRVAttrUtils.inc"
} // namespace spirv
} // namespace mlir
//===----------------------------------------------------------------------===//
// DictionaryDict derived attributes
//===----------------------------------------------------------------------===//
#include "mlir/Dialect/SPIRV/IR/TargetAndABI.cpp.inc"
namespace mlir {
//===----------------------------------------------------------------------===//
// Attribute storage classes
@ -292,7 +286,7 @@ spirv::TargetEnvAttr spirv::TargetEnvAttr::get(spirv::VerCapExtAttr triple,
Vendor vendorID,
DeviceType deviceType,
uint32_t deviceID,
DictionaryAttr limits) {
ResourceLimitsAttr limits) {
assert(triple && limits && "expected valid triple and limits");
MLIRContext *context = triple.getContext();
return Base::get(context, triple, vendorID, deviceType, deviceID, limits);
@ -340,16 +334,332 @@ spirv::ResourceLimitsAttr spirv::TargetEnvAttr::getResourceLimits() const {
return getImpl()->limits.cast<spirv::ResourceLimitsAttr>();
}
LogicalResult
spirv::TargetEnvAttr::verify(function_ref<InFlightDiagnostic()> emitError,
spirv::VerCapExtAttr /*triple*/,
spirv::Vendor /*vendorID*/,
spirv::DeviceType /*deviceType*/,
uint32_t /*deviceID*/, DictionaryAttr limits) {
if (!limits.isa<spirv::ResourceLimitsAttr>())
return emitError() << "expected spirv::ResourceLimitsAttr for limits";
//===----------------------------------------------------------------------===//
// ODS Generated Attributes
//===----------------------------------------------------------------------===//
return success();
#define GET_ATTRDEF_CLASSES
#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.cpp.inc"
//===----------------------------------------------------------------------===//
// Attribute Parsing
//===----------------------------------------------------------------------===//
/// Parses a comma-separated list of keywords, invokes `processKeyword` on each
/// of the parsed keyword, and returns failure if any error occurs.
static ParseResult
parseKeywordList(DialectAsmParser &parser,
function_ref<LogicalResult(SMLoc, StringRef)> processKeyword) {
if (parser.parseLSquare())
return failure();
// Special case for empty list.
if (succeeded(parser.parseOptionalRSquare()))
return success();
// Keep parsing the keyword and an optional comma following it. If the comma
// is successfully parsed, then we have more keywords to parse.
if (failed(parser.parseCommaSeparatedList([&]() {
auto loc = parser.getCurrentLocation();
StringRef keyword;
if (parser.parseKeyword(&keyword) ||
failed(processKeyword(loc, keyword)))
return failure();
return success();
})))
return failure();
return parser.parseRSquare();
}
/// Parses a spirv::InterfaceVarABIAttr.
static Attribute parseInterfaceVarABIAttr(DialectAsmParser &parser) {
if (parser.parseLess())
return {};
Builder &builder = parser.getBuilder();
if (parser.parseLParen())
return {};
IntegerAttr descriptorSetAttr;
{
auto loc = parser.getCurrentLocation();
uint32_t descriptorSet = 0;
auto descriptorSetParseResult = parser.parseOptionalInteger(descriptorSet);
if (!descriptorSetParseResult.hasValue() ||
failed(*descriptorSetParseResult)) {
parser.emitError(loc, "missing descriptor set");
return {};
}
descriptorSetAttr = builder.getI32IntegerAttr(descriptorSet);
}
if (parser.parseComma())
return {};
IntegerAttr bindingAttr;
{
auto loc = parser.getCurrentLocation();
uint32_t binding = 0;
auto bindingParseResult = parser.parseOptionalInteger(binding);
if (!bindingParseResult.hasValue() || failed(*bindingParseResult)) {
parser.emitError(loc, "missing binding");
return {};
}
bindingAttr = builder.getI32IntegerAttr(binding);
}
if (parser.parseRParen())
return {};
IntegerAttr storageClassAttr;
{
if (succeeded(parser.parseOptionalComma())) {
auto loc = parser.getCurrentLocation();
StringRef storageClass;
if (parser.parseKeyword(&storageClass))
return {};
if (auto storageClassSymbol =
spirv::symbolizeStorageClass(storageClass)) {
storageClassAttr = builder.getI32IntegerAttr(
static_cast<uint32_t>(*storageClassSymbol));
} else {
parser.emitError(loc, "unknown storage class: ") << storageClass;
return {};
}
}
}
if (parser.parseGreater())
return {};
return spirv::InterfaceVarABIAttr::get(descriptorSetAttr, bindingAttr,
storageClassAttr);
}
static Attribute parseVerCapExtAttr(DialectAsmParser &parser) {
if (parser.parseLess())
return {};
Builder &builder = parser.getBuilder();
IntegerAttr versionAttr;
{
auto loc = parser.getCurrentLocation();
StringRef version;
if (parser.parseKeyword(&version) || parser.parseComma())
return {};
if (auto versionSymbol = spirv::symbolizeVersion(version)) {
versionAttr =
builder.getI32IntegerAttr(static_cast<uint32_t>(*versionSymbol));
} else {
parser.emitError(loc, "unknown version: ") << version;
return {};
}
}
ArrayAttr capabilitiesAttr;
{
SmallVector<Attribute, 4> capabilities;
SMLoc errorloc;
StringRef errorKeyword;
auto processCapability = [&](SMLoc loc, StringRef capability) {
if (auto capSymbol = spirv::symbolizeCapability(capability)) {
capabilities.push_back(
builder.getI32IntegerAttr(static_cast<uint32_t>(*capSymbol)));
return success();
}
return errorloc = loc, errorKeyword = capability, failure();
};
if (parseKeywordList(parser, processCapability) || parser.parseComma()) {
if (!errorKeyword.empty())
parser.emitError(errorloc, "unknown capability: ") << errorKeyword;
return {};
}
capabilitiesAttr = builder.getArrayAttr(capabilities);
}
ArrayAttr extensionsAttr;
{
SmallVector<Attribute, 1> extensions;
SMLoc errorloc;
StringRef errorKeyword;
auto processExtension = [&](SMLoc loc, StringRef extension) {
if (spirv::symbolizeExtension(extension)) {
extensions.push_back(builder.getStringAttr(extension));
return success();
}
return errorloc = loc, errorKeyword = extension, failure();
};
if (parseKeywordList(parser, processExtension)) {
if (!errorKeyword.empty())
parser.emitError(errorloc, "unknown extension: ") << errorKeyword;
return {};
}
extensionsAttr = builder.getArrayAttr(extensions);
}
if (parser.parseGreater())
return {};
return spirv::VerCapExtAttr::get(versionAttr, capabilitiesAttr,
extensionsAttr);
}
/// Parses a spirv::TargetEnvAttr.
static Attribute parseTargetEnvAttr(DialectAsmParser &parser) {
if (parser.parseLess())
return {};
spirv::VerCapExtAttr tripleAttr;
if (parser.parseAttribute(tripleAttr) || parser.parseComma())
return {};
// Parse [vendor[:device-type[:device-id]]]
Vendor vendorID = Vendor::Unknown;
DeviceType deviceType = DeviceType::Unknown;
uint32_t deviceID = spirv::TargetEnvAttr::kUnknownDeviceID;
{
auto loc = parser.getCurrentLocation();
StringRef vendorStr;
if (succeeded(parser.parseOptionalKeyword(&vendorStr))) {
if (auto vendorSymbol = spirv::symbolizeVendor(vendorStr)) {
vendorID = *vendorSymbol;
} else {
parser.emitError(loc, "unknown vendor: ") << vendorStr;
}
if (succeeded(parser.parseOptionalColon())) {
loc = parser.getCurrentLocation();
StringRef deviceTypeStr;
if (parser.parseKeyword(&deviceTypeStr))
return {};
if (auto deviceTypeSymbol = spirv::symbolizeDeviceType(deviceTypeStr)) {
deviceType = *deviceTypeSymbol;
} else {
parser.emitError(loc, "unknown device type: ") << deviceTypeStr;
}
if (succeeded(parser.parseOptionalColon())) {
loc = parser.getCurrentLocation();
if (parser.parseInteger(deviceID))
return {};
}
}
if (parser.parseComma())
return {};
}
}
ResourceLimitsAttr limitsAttr;
if (parser.parseAttribute(limitsAttr) || parser.parseGreater())
return {};
return spirv::TargetEnvAttr::get(tripleAttr, vendorID, deviceType, deviceID,
limitsAttr);
}
Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser,
Type type) const {
// SPIR-V attributes are dictionaries so they do not have type.
if (type) {
parser.emitError(parser.getNameLoc(), "unexpected type");
return {};
}
// Parse the kind keyword first.
StringRef attrKind;
if (parser.parseKeyword(&attrKind))
return {};
Attribute attr;
OptionalParseResult result =
generatedAttributeParser(parser, attrKind, type, attr);
if (result.hasValue()) {
if (failed(result.getValue()))
return {};
return attr;
}
if (attrKind == spirv::TargetEnvAttr::getKindName())
return parseTargetEnvAttr(parser);
if (attrKind == spirv::VerCapExtAttr::getKindName())
return parseVerCapExtAttr(parser);
if (attrKind == spirv::InterfaceVarABIAttr::getKindName())
return parseInterfaceVarABIAttr(parser);
parser.emitError(parser.getNameLoc(), "unknown SPIR-V attribute kind: ")
<< attrKind;
return {};
}
//===----------------------------------------------------------------------===//
// Attribute Printing
//===----------------------------------------------------------------------===//
static void print(spirv::VerCapExtAttr triple, DialectAsmPrinter &printer) {
auto &os = printer.getStream();
printer << spirv::VerCapExtAttr::getKindName() << "<"
<< spirv::stringifyVersion(triple.getVersion()) << ", [";
llvm::interleaveComma(
triple.getCapabilities(), os,
[&](spirv::Capability cap) { os << spirv::stringifyCapability(cap); });
printer << "], [";
llvm::interleaveComma(triple.getExtensionsAttr(), os, [&](Attribute attr) {
os << attr.cast<StringAttr>().getValue();
});
printer << "]>";
}
static void print(spirv::TargetEnvAttr targetEnv, DialectAsmPrinter &printer) {
printer << spirv::TargetEnvAttr::getKindName() << "<#spv.";
print(targetEnv.getTripleAttr(), printer);
spirv::Vendor vendorID = targetEnv.getVendorID();
spirv::DeviceType deviceType = targetEnv.getDeviceType();
uint32_t deviceID = targetEnv.getDeviceID();
if (vendorID != spirv::Vendor::Unknown) {
printer << ", " << spirv::stringifyVendor(vendorID);
if (deviceType != spirv::DeviceType::Unknown) {
printer << ":" << spirv::stringifyDeviceType(deviceType);
if (deviceID != spirv::TargetEnvAttr::kUnknownDeviceID)
printer << ":" << deviceID;
}
}
printer << ", " << targetEnv.getResourceLimits() << ">";
}
static void print(spirv::InterfaceVarABIAttr interfaceVarABIAttr,
DialectAsmPrinter &printer) {
printer << spirv::InterfaceVarABIAttr::getKindName() << "<("
<< interfaceVarABIAttr.getDescriptorSet() << ", "
<< interfaceVarABIAttr.getBinding() << ")";
auto storageClass = interfaceVarABIAttr.getStorageClass();
if (storageClass)
printer << ", " << spirv::stringifyStorageClass(*storageClass);
printer << ">";
}
void SPIRVDialect::printAttribute(Attribute attr,
DialectAsmPrinter &printer) const {
if (succeeded(generatedAttributePrinter(attr, printer)))
return;
if (auto targetEnv = attr.dyn_cast<TargetEnvAttr>())
print(targetEnv, printer);
else if (auto vceAttr = attr.dyn_cast<VerCapExtAttr>())
print(vceAttr, printer);
else if (auto interfaceVarABIAttr = attr.dyn_cast<InterfaceVarABIAttr>())
print(interfaceVarABIAttr, printer);
else
llvm_unreachable("unhandled SPIR-V attribute kind");
}
//===----------------------------------------------------------------------===//
@ -358,4 +668,8 @@ spirv::TargetEnvAttr::verify(function_ref<InFlightDiagnostic()> emitError,
void spirv::SPIRVDialect::registerAttributes() {
addAttributes<InterfaceVarABIAttr, TargetEnvAttr, VerCapExtAttr>();
addAttributes<
#define GET_ATTRDEF_LIST
#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.cpp.inc"
>();
}

333
mlir/lib/Dialect/SPIRV/IR/SPIRVDialect.cpp
View File

@ -870,330 +870,6 @@ void SPIRVDialect::printType(Type type, DialectAsmPrinter &os) const {
.Default([](Type) { llvm_unreachable("unhandled SPIR-V type"); });
}
//===----------------------------------------------------------------------===//
// Attribute Parsing
//===----------------------------------------------------------------------===//
/// Parses a comma-separated list of keywords, invokes `processKeyword` on each
/// of the parsed keyword, and returns failure if any error occurs.
static ParseResult parseKeywordList(
DialectAsmParser &parser,
function_ref<LogicalResult(SMLoc, StringRef)> processKeyword) {
if (parser.parseLSquare())
return failure();
// Special case for empty list.
if (succeeded(parser.parseOptionalRSquare()))
return success();
// Keep parsing the keyword and an optional comma following it. If the comma
// is successfully parsed, then we have more keywords to parse.
if (failed(parser.parseCommaSeparatedList([&]() {
auto loc = parser.getCurrentLocation();
StringRef keyword;
if (parser.parseKeyword(&keyword) ||
failed(processKeyword(loc, keyword)))
return failure();
return success();
})))
return failure();
return parser.parseRSquare();
}
/// Parses a spirv::InterfaceVarABIAttr.
static Attribute parseInterfaceVarABIAttr(DialectAsmParser &parser) {
if (parser.parseLess())
return {};
Builder &builder = parser.getBuilder();
if (parser.parseLParen())
return {};
IntegerAttr descriptorSetAttr;
{
auto loc = parser.getCurrentLocation();
uint32_t descriptorSet = 0;
auto descriptorSetParseResult = parser.parseOptionalInteger(descriptorSet);
if (!descriptorSetParseResult.hasValue() ||
failed(*descriptorSetParseResult)) {
parser.emitError(loc, "missing descriptor set");
return {};
}
descriptorSetAttr = builder.getI32IntegerAttr(descriptorSet);
}
if (parser.parseComma())
return {};
IntegerAttr bindingAttr;
{
auto loc = parser.getCurrentLocation();
uint32_t binding = 0;
auto bindingParseResult = parser.parseOptionalInteger(binding);
if (!bindingParseResult.hasValue() || failed(*bindingParseResult)) {
parser.emitError(loc, "missing binding");
return {};
}
bindingAttr = builder.getI32IntegerAttr(binding);
}
if (parser.parseRParen())
return {};
IntegerAttr storageClassAttr;
{
if (succeeded(parser.parseOptionalComma())) {
auto loc = parser.getCurrentLocation();
StringRef storageClass;
if (parser.parseKeyword(&storageClass))
return {};
if (auto storageClassSymbol =
spirv::symbolizeStorageClass(storageClass)) {
storageClassAttr = builder.getI32IntegerAttr(
static_cast<uint32_t>(*storageClassSymbol));
} else {
parser.emitError(loc, "unknown storage class: ") << storageClass;
return {};
}
}
}
if (parser.parseGreater())
return {};
return spirv::InterfaceVarABIAttr::get(descriptorSetAttr, bindingAttr,
storageClassAttr);
}
static Attribute parseVerCapExtAttr(DialectAsmParser &parser) {
if (parser.parseLess())
return {};
Builder &builder = parser.getBuilder();
IntegerAttr versionAttr;
{
auto loc = parser.getCurrentLocation();
StringRef version;
if (parser.parseKeyword(&version) || parser.parseComma())
return {};
if (auto versionSymbol = spirv::symbolizeVersion(version)) {
versionAttr =
builder.getI32IntegerAttr(static_cast<uint32_t>(*versionSymbol));
} else {
parser.emitError(loc, "unknown version: ") << version;
return {};
}
}
ArrayAttr capabilitiesAttr;
{
SmallVector<Attribute, 4> capabilities;
SMLoc errorloc;
StringRef errorKeyword;
auto processCapability = [&](SMLoc loc, StringRef capability) {
if (auto capSymbol = spirv::symbolizeCapability(capability)) {
capabilities.push_back(
builder.getI32IntegerAttr(static_cast<uint32_t>(*capSymbol)));
return success();
}
return errorloc = loc, errorKeyword = capability, failure();
};
if (parseKeywordList(parser, processCapability) || parser.parseComma()) {
if (!errorKeyword.empty())
parser.emitError(errorloc, "unknown capability: ") << errorKeyword;
return {};
}
capabilitiesAttr = builder.getArrayAttr(capabilities);
}
ArrayAttr extensionsAttr;
{
SmallVector<Attribute, 1> extensions;
SMLoc errorloc;
StringRef errorKeyword;
auto processExtension = [&](SMLoc loc, StringRef extension) {
if (spirv::symbolizeExtension(extension)) {
extensions.push_back(builder.getStringAttr(extension));
return success();
}
return errorloc = loc, errorKeyword = extension, failure();
};
if (parseKeywordList(parser, processExtension)) {
if (!errorKeyword.empty())
parser.emitError(errorloc, "unknown extension: ") << errorKeyword;
return {};
}
extensionsAttr = builder.getArrayAttr(extensions);
}
if (parser.parseGreater())
return {};
return spirv::VerCapExtAttr::get(versionAttr, capabilitiesAttr,
extensionsAttr);
}
/// Parses a spirv::TargetEnvAttr.
static Attribute parseTargetEnvAttr(DialectAsmParser &parser) {
if (parser.parseLess())
return {};
spirv::VerCapExtAttr tripleAttr;
if (parser.parseAttribute(tripleAttr) || parser.parseComma())
return {};
// Parse [vendor[:device-type[:device-id]]]
Vendor vendorID = Vendor::Unknown;
DeviceType deviceType = DeviceType::Unknown;
uint32_t deviceID = spirv::TargetEnvAttr::kUnknownDeviceID;
{
auto loc = parser.getCurrentLocation();
StringRef vendorStr;
if (succeeded(parser.parseOptionalKeyword(&vendorStr))) {
if (auto vendorSymbol = spirv::symbolizeVendor(vendorStr)) {
vendorID = *vendorSymbol;
} else {
parser.emitError(loc, "unknown vendor: ") << vendorStr;
}
if (succeeded(parser.parseOptionalColon())) {
loc = parser.getCurrentLocation();
StringRef deviceTypeStr;
if (parser.parseKeyword(&deviceTypeStr))
return {};
if (auto deviceTypeSymbol = spirv::symbolizeDeviceType(deviceTypeStr)) {
deviceType = *deviceTypeSymbol;
} else {
parser.emitError(loc, "unknown device type: ") << deviceTypeStr;
}
if (succeeded(parser.parseOptionalColon())) {
loc = parser.getCurrentLocation();
if (parser.parseInteger(deviceID))
return {};
}
}
if (parser.parseComma())
return {};
}
}
DictionaryAttr limitsAttr;
{
auto loc = parser.getCurrentLocation();
if (parser.parseAttribute(limitsAttr))
return {};
if (!limitsAttr.isa<spirv::ResourceLimitsAttr>()) {
parser.emitError(
loc,
"limits must be a dictionary attribute containing two 32-bit integer "
"attributes 'max_compute_workgroup_invocations' and "
"'max_compute_workgroup_size'");
return {};
}
}
if (parser.parseGreater())
return {};
return spirv::TargetEnvAttr::get(tripleAttr, vendorID, deviceType, deviceID,
limitsAttr);
}
Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser,
Type type) const {
// SPIR-V attributes are dictionaries so they do not have type.
if (type) {
parser.emitError(parser.getNameLoc(), "unexpected type");
return {};
}
// Parse the kind keyword first.
StringRef attrKind;
if (parser.parseKeyword(&attrKind))
return {};
if (attrKind == spirv::TargetEnvAttr::getKindName())
return parseTargetEnvAttr(parser);
if (attrKind == spirv::VerCapExtAttr::getKindName())
return parseVerCapExtAttr(parser);
if (attrKind == spirv::InterfaceVarABIAttr::getKindName())
return parseInterfaceVarABIAttr(parser);
parser.emitError(parser.getNameLoc(), "unknown SPIR-V attribute kind: ")
<< attrKind;
return {};
}
//===----------------------------------------------------------------------===//
// Attribute Printing
//===----------------------------------------------------------------------===//
static void print(spirv::VerCapExtAttr triple, DialectAsmPrinter &printer) {
auto &os = printer.getStream();
printer << spirv::VerCapExtAttr::getKindName() << "<"
<< spirv::stringifyVersion(triple.getVersion()) << ", [";
llvm::interleaveComma(
triple.getCapabilities(), os,
[&](spirv::Capability cap) { os << spirv::stringifyCapability(cap); });
printer << "], [";
llvm::interleaveComma(triple.getExtensionsAttr(), os, [&](Attribute attr) {
os << attr.cast<StringAttr>().getValue();
});
printer << "]>";
}
static void print(spirv::TargetEnvAttr targetEnv, DialectAsmPrinter &printer) {
printer << spirv::TargetEnvAttr::getKindName() << "<#spv.";
print(targetEnv.getTripleAttr(), printer);
spirv::Vendor vendorID = targetEnv.getVendorID();
spirv::DeviceType deviceType = targetEnv.getDeviceType();
uint32_t deviceID = targetEnv.getDeviceID();
if (vendorID != spirv::Vendor::Unknown) {
printer << ", " << spirv::stringifyVendor(vendorID);
if (deviceType != spirv::DeviceType::Unknown) {
printer << ":" << spirv::stringifyDeviceType(deviceType);
if (deviceID != spirv::TargetEnvAttr::kUnknownDeviceID)
printer << ":" << deviceID;
}
}
printer << ", " << targetEnv.getResourceLimits() << ">";
}
static void print(spirv::InterfaceVarABIAttr interfaceVarABIAttr,
DialectAsmPrinter &printer) {
printer << spirv::InterfaceVarABIAttr::getKindName() << "<("
<< interfaceVarABIAttr.getDescriptorSet() << ", "
<< interfaceVarABIAttr.getBinding() << ")";
auto storageClass = interfaceVarABIAttr.getStorageClass();
if (storageClass)
printer << ", " << spirv::stringifyStorageClass(*storageClass);
printer << ">";
}
void SPIRVDialect::printAttribute(Attribute attr,
DialectAsmPrinter &printer) const {
if (auto targetEnv = attr.dyn_cast<TargetEnvAttr>())
print(targetEnv, printer);
else if (auto vceAttr = attr.dyn_cast<VerCapExtAttr>())
print(vceAttr, printer);
else if (auto interfaceVarABIAttr = attr.dyn_cast<InterfaceVarABIAttr>())
print(interfaceVarABIAttr, printer);
else
llvm_unreachable("unhandled SPIR-V attribute kind");
}
//===----------------------------------------------------------------------===//
// Constant
//===----------------------------------------------------------------------===//
@ -1216,14 +892,11 @@ LogicalResult SPIRVDialect::verifyOperationAttribute(Operation *op,
StringRef symbol = attribute.getName().strref();
Attribute attr = attribute.getValue();
// TODO: figure out a way to generate the description from the
// StructAttr definition.
if (symbol == spirv::getEntryPointABIAttrName()) {
if (!attr.isa<spirv::EntryPointABIAttr>())
if (!attr.isa<spirv::EntryPointABIAttr>()) {
return op->emitError("'")
<< symbol
<< "' attribute must be a dictionary attribute containing one "
"32-bit integer elements attribute: 'local_size'";
<< symbol << "' attribute must be an entry point ABI attribute";
}
} else if (symbol == spirv::getTargetEnvAttrName()) {
if (!attr.isa<spirv::TargetEnvAttr>())
return op->emitError("'") << symbol << "' must be a spirv::TargetEnvAttr";

25
mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp
View File

@ -121,14 +121,13 @@ StringRef spirv::getEntryPointABIAttrName() { return "spv.entry_point_abi"; }
spirv::EntryPointABIAttr
spirv::getEntryPointABIAttr(ArrayRef<int32_t> localSize, MLIRContext *context) {
if (localSize.empty())
return spirv::EntryPointABIAttr::get(nullptr, context);
return spirv::EntryPointABIAttr::get(context, nullptr);
assert(localSize.size() == 3);
return spirv::EntryPointABIAttr::get(
DenseElementsAttr::get<int32_t>(
VectorType::get(3, IntegerType::get(context, 32)), localSize)
.cast<DenseIntElementsAttr>(),
context);
context, DenseElementsAttr::get<int32_t>(
VectorType::get(3, IntegerType::get(context, 32)), localSize)
.cast<DenseIntElementsAttr>());
}
spirv::EntryPointABIAttr spirv::lookupEntryPointABI(Operation *op) {
@ -146,7 +145,7 @@ spirv::EntryPointABIAttr spirv::lookupEntryPointABI(Operation *op) {
DenseIntElementsAttr spirv::lookupLocalWorkGroupSize(Operation *op) {
if (auto entryPoint = spirv::lookupEntryPointABI(op))
return entryPoint.local_size();
return entryPoint.getLocal_size();
return {};
}
@ -155,12 +154,14 @@ spirv::ResourceLimitsAttr
spirv::getDefaultResourceLimits(MLIRContext *context) {
// All the fields have default values. Here we just provide a nicer way to
// construct a default resource limit attribute.
return spirv::ResourceLimitsAttr ::get(
/*max_compute_shared_memory_size=*/nullptr,
/*max_compute_workgroup_invocations=*/nullptr,
/*max_compute_workgroup_size=*/nullptr,
/*subgroup_size=*/nullptr,
/*cooperative_matrix_properties_nv=*/nullptr, context);
Builder b(context);
return spirv::ResourceLimitsAttr::get(
context,
/*max_compute_shared_memory_size=*/16384,
/*max_compute_workgroup_invocations=*/128,
/*max_compute_workgroup_size=*/b.getI32ArrayAttr({128, 128, 64}),
/*subgroup_size=*/32,
/*cooperative_matrix_properties_nv=*/ArrayAttr());
}
StringRef spirv::getTargetEnvAttrName() { return "spv.target_env"; }

2
mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp
View File

@ -135,7 +135,7 @@ static LogicalResult lowerEntryPointABIAttr(spirv::FuncOp funcOp,
funcOp.getLoc(), executionModel.getValue(), funcOp, interfaceVars);
// Specifies the spv.ExecutionModeOp.
auto localSizeAttr = entryPointAttr.local_size();
auto localSizeAttr = entryPointAttr.getLocal_size();
if (localSizeAttr) {
auto values = localSizeAttr.getValues<int32_t>();
SmallVector<int32_t, 3> localSize(values);

56
mlir/test/Conversion/ArithmeticToSPIRV/arithmetic-to-spirv.mlir
View File

@ -6,7 +6,7 @@
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64, Shader], []>, {}>
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64, Shader], []>, #spv.resource_limits<>>
} {
// Check integer operation conversions.
@ -154,7 +154,7 @@ func.func @unsupported_2x2elem_vector(%arg0: vector<2x2xi32>) {
// Check that types are converted to 32-bit when no special capabilities.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @int_vector23
@ -182,7 +182,7 @@ func.func @float_scalar(%arg0: f16, %arg1: f64) {
// Check that types are converted to 32-bit when no special capabilities that
// are not supported.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// expected-error @+1 {{failed to materialize conversion for block argument #0 that remained live after conversion, type was 'vector<4xi64>', with target type 'vector<4xi32>'}}
@ -202,7 +202,7 @@ func.func @int_vector4_invalid(%arg0: vector<4xi64>) {
//===----------------------------------------------------------------------===//
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @bitwise_scalar
@ -280,7 +280,7 @@ func.func @shift_vector(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) {
//===----------------------------------------------------------------------===//
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @cmpf
@ -318,7 +318,7 @@ func.func @cmpf(%arg0 : f32, %arg1 : f32) {
// With Kernel capability, we can convert NaN check to spv.Ordered/spv.Unordered.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @cmpf
@ -336,7 +336,7 @@ func.func @cmpf(%arg0 : f32, %arg1 : f32) {
// Without Kernel capability, we need to convert NaN check to spv.IsNan.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @cmpf
@ -364,7 +364,7 @@ func.func @cmpf(%arg0 : f32, %arg1 : f32) {
//===----------------------------------------------------------------------===//
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @cmpi
@ -420,7 +420,7 @@ func.func @vecboolcmpi(%arg0 : vector<4xi1>, %arg1 : vector<4xi1>) {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @constant
@ -493,7 +493,7 @@ func.func @constant_size1() {
// Check that constants are converted to 32-bit when no special capability.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @constant_16bit
@ -585,7 +585,7 @@ func.func @unsupported_cases() {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: index_cast1
@ -833,7 +833,7 @@ func.func @fptosi2(%arg0 : f16) -> i16 {
// Checks that cast types will be adjusted when missing special capabilities for
// certain non-32-bit scalar types.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float64], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float64], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @fpext1
@ -859,7 +859,7 @@ func.func @fpext2(%arg0 : f32) -> f64 {
// Checks that cast types will be adjusted when missing special capabilities for
// certain non-32-bit scalar types.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @fptrunc1
@ -892,7 +892,7 @@ func.func @sitofp(%arg0 : i64) -> f64 {
// Check OpenCL lowering of arith.remsi
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Int16, Kernel], []>, {}>
#spv.vce<v1.0, [Int16, Kernel], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @scalar_srem
@ -928,7 +928,7 @@ func.func @vector_srem(%arg0: vector<3xi16>, %arg1: vector<3xi16>) {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader, Int8, Int16, Int64, Float16, Float64],
[SPV_KHR_storage_buffer_storage_class]>, {}>
[SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @select
@ -949,7 +949,7 @@ func.func @select(%arg0 : i32, %arg1 : i32) {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64, Shader], []>, {}>
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64, Shader], []>, #spv.resource_limits<>>
} {
// Check integer operation conversions.
@ -1079,7 +1079,7 @@ func.func @unsupported_2x2elem_vector(%arg0: vector<2x2xi32>) {
// Check that types are converted to 32-bit when no special capabilities.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @int_vector23
@ -1107,7 +1107,7 @@ func.func @float_scalar(%arg0: f16, %arg1: f64) {
// Check that types are converted to 32-bit when no special capabilities that
// are not supported.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// expected-error@below {{failed to materialize conversion for block argument #0 that remained live after conversion}}
@ -1127,7 +1127,7 @@ func.func @int_vector4_invalid(%arg0: vector<4xi64>) {
//===----------------------------------------------------------------------===//
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @bitwise_scalar
@ -1205,7 +1205,7 @@ func.func @shift_vector(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) {
//===----------------------------------------------------------------------===//
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @cmpf
@ -1243,7 +1243,7 @@ func.func @cmpf(%arg0 : f32, %arg1 : f32) {
// With Kernel capability, we can convert NaN check to spv.Ordered/spv.Unordered.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @cmpf
@ -1261,7 +1261,7 @@ func.func @cmpf(%arg0 : f32, %arg1 : f32) {
// Without Kernel capability, we need to convert NaN check to spv.IsNan.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @cmpf
@ -1289,7 +1289,7 @@ func.func @cmpf(%arg0 : f32, %arg1 : f32) {
//===----------------------------------------------------------------------===//
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @cmpi
@ -1345,7 +1345,7 @@ func.func @vecboolcmpi(%arg0 : vector<4xi1>, %arg1 : vector<4xi1>) {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @constant
@ -1407,7 +1407,7 @@ func.func @constant_64bit() {
// Check that constants are converted to 32-bit when no special capability.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @constant_16bit
@ -1490,7 +1490,7 @@ func.func @unsupported_cases() {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: index_cast1
@ -1729,7 +1729,7 @@ func.func @fptosi2(%arg0 : f16) -> i16 {
// Checks that cast types will be adjusted when missing special capabilities for
// certain non-32-bit scalar types.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float64], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float64], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @fpext1
@ -1755,7 +1755,7 @@ func.func @fpext2(%arg0 : f32) -> f64 {
// Checks that cast types will be adjusted when missing special capabilities for
// certain non-32-bit scalar types.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @fptrunc1

2
mlir/test/Conversion/ControlFlowToSPIRV/cf-ops-to-spirv.mlir
View File

@ -5,7 +5,7 @@
//===----------------------------------------------------------------------===//
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: func @simple_loop

2
mlir/test/Conversion/FuncToSPIRV/func-ops-to-spirv.mlir
View File

@ -5,7 +5,7 @@
//===----------------------------------------------------------------------===//
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @return_none_val

46
mlir/test/Conversion/FuncToSPIRV/types-to-spirv.mlir
View File

@ -8,7 +8,7 @@
// Check that non-32-bit integer types are converted to 32-bit types if the
// corresponding capabilities are not available.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @integer8
@ -48,7 +48,7 @@ func.func @integer64(%arg0: i64, %arg1: si64, %arg2: ui64) { return }
// Check that non-32-bit integer types are kept untouched if the corresponding
// capabilities are available.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Int8, Int16, Int64], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Int8, Int16, Int64], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @integer8
@ -87,7 +87,7 @@ func.func @integer64(%arg0: i64, %arg1: si64, %arg2: ui64) { return }
// Check that weird bitwidths are not supported.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-NOT: spv.func @integer4
@ -108,7 +108,7 @@ func.func @integer42(%arg0: i42) { return }
// The index type is always converted into i32.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @index_type
@ -126,7 +126,7 @@ func.func @index_type(%arg0: index) { return }
// Check that non-32-bit float types are converted to 32-bit types if the
// corresponding capabilities are not available.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @float16
@ -148,7 +148,7 @@ func.func @float64(%arg0: f64) { return }
// Check that non-32-bit float types are kept untouched if the corresponding
// capabilities are available.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16, Float64], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16, Float64], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @float16
@ -169,7 +169,7 @@ func.func @float64(%arg0: f64) { return }
// Check that bf16 is not supported.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-NOT: spv.func @bf16_type
@ -186,7 +186,7 @@ func.func @bf16_type(%arg0: bf16) { return }
// Check that capabilities for scalar types affects vector types too: no special
// capabilities available means using turning element types to 32-bit.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @int_vector
@ -215,7 +215,7 @@ func.func @float_vector(
// special capabilities means keep vector types untouched.
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @int_vector
@ -246,7 +246,7 @@ func.func @one_element_vector(%arg0: vector<1xi32>) { return }
// Check that > 4-element vectors are not supported.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-NOT: spv.func @large_vector
@ -263,7 +263,7 @@ func.func @large_vector(%arg0: vector<1024xi32>) { return }
// Check memory spaces.
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: func @memref_mem_space
@ -301,7 +301,7 @@ func.func @memref_1bit_type(
// requires special capability and extension: convert them to 32-bit if not
// satisfied.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// An i1 is store in 8-bit, so 5xi1 has 40 bits, which is stored in 2xi32.
@ -399,7 +399,7 @@ func.func @memref_64bit_Output(%arg4: memref<16xf64, 10>) { return }
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [StoragePushConstant8, StoragePushConstant16, Int64, Float64],
[SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, {}>
[SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @memref_8bit_PushConstant
@ -440,7 +440,7 @@ func.func @memref_64bit_PushConstant(
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [StorageBuffer8BitAccess, StorageBuffer16BitAccess, Int64, Float64],
[SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, {}>
[SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @memref_8bit_StorageBuffer
@ -481,7 +481,7 @@ func.func @memref_64bit_StorageBuffer(
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [UniformAndStorageBuffer8BitAccess, StorageUniform16, Int64, Float64],
[SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, {}>
[SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @memref_8bit_Uniform
@ -521,7 +521,7 @@ func.func @memref_64bit_Uniform(
// and extension is available.
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [StorageInputOutput16, Int64, Float64], [SPV_KHR_16bit_storage]>, {}>
#spv.vce<v1.0, [StorageInputOutput16, Int64, Float64], [SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @memref_16bit_Input
@ -565,7 +565,7 @@ func.func @memref_64bit_Output(
// Check that memref offset and strides affect the array size.
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [StorageBuffer16BitAccess], [SPV_KHR_16bit_storage]>, {}>
#spv.vce<v1.0, [StorageBuffer16BitAccess], [SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @memref_offset_strides
@ -599,7 +599,7 @@ func.func @memref_offset_strides(
// Dynamic shapes
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// Check that unranked shapes are not supported.
@ -677,7 +677,7 @@ func.func @memref_16bit_Output(%arg4: memref<?xf16, 10>) { return }
// Vector types
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: func @memref_vector
@ -701,7 +701,7 @@ func.func @dynamic_dim_memref_vector(%arg0: memref<8x?xvector<4xi32>>,
// Vector types, check that sizes not available in SPIR-V are not transformed.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: func @memref_vector_wrong_size
@ -721,7 +721,7 @@ func.func @memref_vector_wrong_size(
// Check that tensor element types are kept untouched with proper capabilities.
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
#spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @int_tensor_types
@ -752,7 +752,7 @@ func.func @float_tensor_types(
// Check that tensor element types are changed to 32-bit without capabilities.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.func @int_tensor_types
@ -783,7 +783,7 @@ func.func @float_tensor_types(
// Check that dynamic shapes are not supported.
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
} {
// CHECK-LABEL: func @unranked_tensor

34
mlir/test/Conversion/GPUToSPIRV/builtins.mlir
View File

@ -12,7 +12,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
gpu.module @kernels {
gpu.func @builtin_workgroup_id_x() kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
// CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
@ -38,7 +38,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
gpu.module @kernels {
gpu.func @builtin_workgroup_id_y() kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
// CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
@ -62,7 +62,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
gpu.module @kernels {
gpu.func @builtin_workgroup_id_z() kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
// CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
@ -85,7 +85,7 @@ module attributes {gpu.container_module} {
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
gpu.module @kernels {
gpu.func @builtin_workgroup_size_x() kernel
attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 1, 1]>: vector<3xi32>>} {
// The constant value is obtained from the spv.entry_point_abi.
// Note that this ignores the workgroup size specification in gpu.launch.
// We may want to define gpu.workgroup_size and convert it to the entry
@ -110,7 +110,7 @@ module attributes {gpu.container_module} {
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
gpu.module @kernels {
gpu.func @builtin_workgroup_size_y() kernel
attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
// The constant value is obtained from the spv.entry_point_abi.
// CHECK: spv.Constant 4 : i32
%0 = gpu.block_dim y
@ -132,7 +132,7 @@ module attributes {gpu.container_module} {
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
gpu.module @kernels {
gpu.func @builtin_workgroup_size_z() kernel
attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
// The constant value is obtained from the spv.entry_point_abi.
// CHECK: spv.Constant 1 : i32
%0 = gpu.block_dim z
@ -155,7 +155,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
gpu.module @kernels {
gpu.func @builtin_local_id_x() kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[LOCALINVOCATIONID]]
// CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
@ -179,7 +179,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
gpu.module @kernels {
gpu.func @builtin_num_workgroups_x() kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[NUMWORKGROUPS]]
// CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
@ -196,7 +196,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId")
gpu.module @kernels {
gpu.func @builtin_subgroup_id() kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[SUBGROUPID]]
// CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
%0 = gpu.subgroup_id : index
@ -212,7 +212,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups")
gpu.module @kernels {
gpu.func @builtin_num_subgroups() kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[NUMSUBGROUPS]]
// CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
%0 = gpu.num_subgroups : index
@ -235,7 +235,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
gpu.module @kernels {
gpu.func @builtin_workgroup_size_x() kernel
attributes {spv.entry_point_abi = {}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPSIZE]]
// CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
@ -259,7 +259,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
gpu.module @kernels {
gpu.func @builtin_workgroup_size_y() kernel
attributes {spv.entry_point_abi = {}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPSIZE]]
// CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
@ -283,7 +283,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
gpu.module @kernels {
gpu.func @builtin_workgroup_size_z() kernel
attributes {spv.entry_point_abi = {}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPSIZE]]
// CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
@ -307,7 +307,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
gpu.module @kernels {
gpu.func @builtin_global_id_x() kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[GLOBALINVOCATIONID]]
// CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
@ -331,7 +331,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
gpu.module @kernels {
gpu.func @builtin_global_id_y() kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[GLOBALINVOCATIONID]]
// CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
@ -355,7 +355,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
gpu.module @kernels {
gpu.func @builtin_global_id_z() kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[GLOBALINVOCATIONID]]
// CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
@ -373,7 +373,7 @@ module attributes {gpu.container_module} {
// CHECK: spv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize")
gpu.module @kernels {
gpu.func @builtin_subgroup_size() kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[SUBGROUPSIZE]]
// CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
%0 = gpu.subgroup_size : index

4
mlir/test/Conversion/GPUToSPIRV/entry-point.mlir
View File

@ -2,10 +2,10 @@
// RUN: mlir-opt -test-spirv-entry-point-abi="workgroup-size=32" %s | FileCheck %s -check-prefix=WG32
// DEFAULT: gpu.func @foo()
// DEFAULT-SAME: spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}
// DEFAULT-SAME: spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<1> : vector<3xi32>>
// WG32: gpu.func @foo()
// WG32-SAME: spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}
// WG32-SAME: spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>
gpu.module @kernels {
gpu.func @foo() kernel {

4
mlir/test/Conversion/GPUToSPIRV/load-store.mlir
View File

@ -3,7 +3,7 @@
module attributes {
gpu.container_module,
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
func.func @load_store(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>) {
%c0 = arith.constant 0 : index
@ -36,7 +36,7 @@ module attributes {
// CHECK-SAME: %[[ARG5:.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 5), StorageBuffer>}
// CHECK-SAME: %[[ARG6:.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 6), StorageBuffer>}
gpu.func @load_store_kernel(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>, %arg3: index, %arg4: index, %arg5: index, %arg6: index) kernel
attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
// CHECK: %[[ADDRESSWORKGROUPID:.*]] = spv.mlir.addressof @[[$WORKGROUPIDVAR]]
// CHECK: %[[WORKGROUPID:.*]] = spv.Load "Input" %[[ADDRESSWORKGROUPID]]
// CHECK: %[[WORKGROUPIDX:.*]] = spv.CompositeExtract %[[WORKGROUPID]]{{\[}}0 : i32{{\]}}

6
mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir
View File

@ -2,7 +2,7 @@
module attributes {
gpu.container_module,
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel, Addresses], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel, Addresses], []>, #spv.resource_limits<>>
} {
gpu.module @kernels {
// CHECK-LABEL: spv.module @{{.*}} Physical64 OpenCL
@ -11,9 +11,9 @@ module attributes {
// CHECK-NOT: spv.interface_var_abi
// CHECK-SAME: {{%.*}}: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, CrossWorkgroup>
// CHECK-NOT: spv.interface_var_abi
// CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}
// CHECK-SAME: spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 4, 1]> : vector<3xi32>>
gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, 11>) kernel
attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
gpu.return
}
}

14
mlir/test/Conversion/GPUToSPIRV/simple.mlir
View File

@ -6,9 +6,9 @@ module attributes {gpu.container_module} {
// CHECK-LABEL: spv.func @basic_module_structure
// CHECK-SAME: {{%.*}}: f32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>}
// CHECK-SAME: {{%.*}}: !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
// CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}
// CHECK-SAME: spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 4, 1]> : vector<3xi32>>
gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32>) kernel
attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
// CHECK: spv.Return
gpu.return
}
@ -35,14 +35,14 @@ module attributes {gpu.container_module} {
// CHECK-SAME: spv.interface_var_abi = #spv.interface_var_abi<(1, 2), StorageBuffer>
// CHECK-SAME: !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
// CHECK-SAME: spv.interface_var_abi = #spv.interface_var_abi<(3, 0)>
// CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}
// CHECK-SAME: spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 4, 1]> : vector<3xi32>>
gpu.func @basic_module_structure_preset_ABI(
%arg0 : f32
{spv.interface_var_abi = #spv.interface_var_abi<(1, 2), StorageBuffer>},
%arg1 : memref<12xf32>
{spv.interface_var_abi = #spv.interface_var_abi<(3, 0)>}) kernel
attributes
{spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
{spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
// CHECK: spv.Return
gpu.return
}
@ -82,7 +82,7 @@ module attributes {gpu.container_module} {
{spv.interface_var_abi = #spv.interface_var_abi<(1, 2), StorageBuffer>},
%arg1 : memref<12xf32>) kernel
attributes
{spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
{spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
gpu.return
}
}
@ -99,7 +99,7 @@ module attributes {gpu.container_module} {
%arg1 : memref<12xf32>
{spv.interface_var_abi = #spv.interface_var_abi<(3, 0)>}) kernel
attributes
{spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
{spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
gpu.return
}
}
@ -111,7 +111,7 @@ module attributes {gpu.container_module} {
gpu.module @kernels {
// CHECK-LABEL: spv.func @barrier
gpu.func @barrier(%arg0 : f32, %arg1 : memref<12xf32>) kernel
attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
// CHECK: spv.ControlBarrier Workgroup, Workgroup, "AcquireRelease|WorkgroupMemory"
gpu.barrier
gpu.return

14
mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir
View File

@ -14,7 +14,7 @@
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
#spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #spv.resource_limits<>>
} {
// CHECK: spv.GlobalVariable
@ -45,7 +45,7 @@ module attributes {
// CHECK: spv.Return
func.func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>) attributes {
spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}
spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>
} {
linalg.generic #single_workgroup_reduction_trait
ins(%input : memref<16xi32>)
@ -72,7 +72,7 @@ func.func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1x
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
#spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #spv.resource_limits<>>
} {
func.func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>) {
// expected-error @+1 {{failed to legalize operation 'linalg.generic'}}
@ -101,10 +101,10 @@ func.func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1x
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
#spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #spv.resource_limits<>>
} {
func.func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>) attributes {
spv.entry_point_abi = {local_size = dense<[32, 1, 1]>: vector<3xi32>}
spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 1, 1]>: vector<3xi32>>
} {
// expected-error @+1 {{failed to legalize operation 'linalg.generic'}}
linalg.generic #single_workgroup_reduction_trait
@ -132,10 +132,10 @@ func.func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1x
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
#spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #spv.resource_limits<>>
} {
func.func @single_workgroup_reduction(%input: memref<16x8xi32>, %output: memref<16xi32>) attributes {
spv.entry_point_abi = {local_size = dense<[16, 8, 1]>: vector<3xi32>}
spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[16, 8, 1]>: vector<3xi32>>
} {
// expected-error @+1 {{failed to legalize operation 'linalg.generic'}}
linalg.generic #single_workgroup_reduction_trait

2
mlir/test/Conversion/MathToSPIRV/math-to-core-spirv.mlir
View File

@ -19,7 +19,7 @@ func.func @copy_sign_scalar(%value: f32, %sign: f32) -> f32 {
// -----
module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16, Int16], []>, {}> } {
module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16, Int16], []>, #spv.resource_limits<>> } {
func.func @copy_sign_vector(%value: vector<3xf16>, %sign: vector<3xf16>) -> vector<3xf16> {
%0 = math.copysign %value, %sign : vector<3xf16>

2
mlir/test/Conversion/MathToSPIRV/math-to-glsl-spirv.mlir
View File

@ -1,6 +1,6 @@
// RUN: mlir-opt -split-input-file -convert-math-to-spirv -verify-diagnostics %s -o - | FileCheck %s
module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], []>, {}> } {
module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], []>, #spv.resource_limits<>> } {
// CHECK-LABEL: @float32_unary_scalar
func.func @float32_unary_scalar(%arg0: f32) {

2
mlir/test/Conversion/MathToSPIRV/math-to-opencl-spirv.mlir
View File

@ -1,6 +1,6 @@
// RUN: mlir-opt -split-input-file -convert-math-to-spirv -verify-diagnostics %s -o - | FileCheck %s
module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel], []>, {}> } {
module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel], []>, #spv.resource_limits<>> } {
// CHECK-LABEL: @float32_unary_scalar
func.func @float32_unary_scalar(%arg0: f32) {

16
mlir/test/Conversion/MemRefToSPIRV/alloc.mlir
View File

@ -2,7 +2,7 @@
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
}
{
func.func @alloc_dealloc_workgroup_mem(%arg0 : index, %arg1 : index) {
@ -27,7 +27,7 @@ module attributes {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
}
{
func.func @alloc_dealloc_workgroup_mem(%arg0 : index, %arg1 : index) {
@ -56,7 +56,7 @@ module attributes {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
}
{
func.func @two_allocs() {
@ -76,7 +76,7 @@ module attributes {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
}
{
func.func @two_allocs_vector() {
@ -97,7 +97,7 @@ module attributes {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
}
{
// CHECK-LABEL: func @alloc_dynamic_size
@ -113,7 +113,7 @@ module attributes {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
}
{
// CHECK-LABEL: func @alloc_unsupported_memory_space
@ -130,7 +130,7 @@ module attributes {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
}
{
// CHECK-LABEL: func @dealloc_dynamic_size
@ -145,7 +145,7 @@ module attributes {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
}
{
// CHECK-LABEL: func @dealloc_unsupported_memory_space

10
mlir/test/Conversion/MemRefToSPIRV/alloca.mlir
View File

@ -1,6 +1,6 @@
// RUN: mlir-opt -split-input-file -convert-memref-to-spirv -canonicalize -verify-diagnostics %s -o - | FileCheck %s
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, #spv.resource_limits<>>} {
func.func @alloc_function_variable(%arg0 : index, %arg1 : index) {
%0 = memref.alloca() : memref<4x5xf32, 6>
%1 = memref.load %0[%arg0, %arg1] : memref<4x5xf32, 6>
@ -19,7 +19,7 @@ module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>
// -----
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, #spv.resource_limits<>>} {
func.func @two_allocs() {
%0 = memref.alloca() : memref<4x5xf32, 6>
%1 = memref.alloca() : memref<2x3xi32, 6>
@ -33,7 +33,7 @@ module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>
// -----
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, #spv.resource_limits<>>} {
func.func @two_allocs_vector() {
%0 = memref.alloca() : memref<4xvector<4xf32>, 6>
%1 = memref.alloca() : memref<2xvector<2xi32>, 6>
@ -48,7 +48,7 @@ module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>
// -----
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, #spv.resource_limits<>>} {
// CHECK-LABEL: func @alloc_dynamic_size
func.func @alloc_dynamic_size(%arg0 : index) -> f32 {
// CHECK: memref.alloca
@ -60,7 +60,7 @@ module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>
// -----
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, #spv.resource_limits<>>} {
// CHECK-LABEL: func @alloc_unsupported_memory_space
func.func @alloc_unsupported_memory_space(%arg0: index) -> f32 {
// CHECK: memref.alloca

6
mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir
View File

@ -11,7 +11,7 @@ module attributes {
StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16,
StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8
],
[SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class]>, {}>
[SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @load_store_zero_rank_float
@ -114,7 +114,7 @@ func.func @store_i1(%dst: memref<4xi1>, %i: index) {
// TODO: Test i64 types.
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @load_i1
@ -314,7 +314,7 @@ func.func @store_f32(%arg0: memref<f32>, %value: f32) {
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Int16, StorageBuffer16BitAccess, Shader],
[SPV_KHR_storage_buffer_storage_class, SPV_KHR_16bit_storage]>, {}>
[SPV_KHR_storage_buffer_storage_class, SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @load_i8

2
mlir/test/Conversion/SCFToSPIRV/for.mlir
View File

@ -2,7 +2,7 @@
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
func.func @loop_kernel(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) {

2
mlir/test/Conversion/SCFToSPIRV/if.mlir
View File

@ -2,7 +2,7 @@
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @kernel_simple_selection

2
mlir/test/Conversion/SCFToSPIRV/while.mlir
View File

@ -2,7 +2,7 @@
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader, Int64], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader, Int64], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: @while_loop1

4
mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
View File

@ -1,6 +1,6 @@
// RUN: mlir-opt --lower-host-to-llvm %s | FileCheck %s
module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, #spv.resource_limits<max_compute_workgroup_invocations = 128, max_compute_workgroup_size = [128, 128, 64]>>} {
// CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() : !llvm.struct<(array<6 x i32>)>
// CHECK: llvm.func @__spv__foo_bar()
@ -32,7 +32,7 @@ module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.v
}
gpu.module @foo {
gpu.func @bar(%arg0: memref<6xi32>) kernel attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}} {
gpu.func @bar(%arg0: memref<6xi32>) kernel attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<1> : vector<3xi32>>} {
gpu.return
}
}

2
mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir
View File

@ -1,6 +1,6 @@
// RUN: mlir-opt -split-input-file -convert-vector-to-spirv -verify-diagnostics %s -o - | FileCheck %s
module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16], []>, {}> } {
module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16], []>, #spv.resource_limits<>> } {
// CHECK-LABEL: @bitcast
// CHECK-SAME: %[[ARG0:.+]]: vector<2xf32>, %[[ARG1:.+]]: vector<2xf16>

89
mlir/test/Dialect/SPIRV/IR/target-and-abi.mlir
View File

@ -26,23 +26,24 @@ func.func @unknown_attr_on_region() -> (i32 {spv.something}) {
// spv.entry_point_abi
//===----------------------------------------------------------------------===//
// expected-error @+1 {{'spv.entry_point_abi' attribute must be a dictionary attribute containing one 32-bit integer elements attribute: 'local_size'}}
// expected-error @+1 {{'spv.entry_point_abi' attribute must be an entry point ABI attribute}}
func.func @spv_entry_point() attributes {
spv.entry_point_abi = 64
} { return }
// -----
// expected-error @+1 {{'spv.entry_point_abi' attribute must be a dictionary attribute containing one 32-bit integer elements attribute: 'local_size'}}
func.func @spv_entry_point() attributes {
spv.entry_point_abi = {local_size = 64}
// expected-error @+2 {{failed to parse SPV_EntryPointABIAttr parameter 'local_size' which is to be a `DenseIntElementsAttr`}}
// expected-error @+1 {{invalid kind of attribute specified}}
spv.entry_point_abi = #spv.entry_point_abi<local_size = 64>
} { return }
// -----
func.func @spv_entry_point() attributes {
// CHECK: {spv.entry_point_abi = {local_size = dense<[64, 1, 1]> : vector<3xi32>}}
spv.entry_point_abi = {local_size = dense<[64, 1, 1]>: vector<3xi32>}
// CHECK: {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[64, 1, 1]> : vector<3xi32>>}
spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[64, 1, 1]>: vector<3xi32>>
} { return }
// -----
@ -104,25 +105,15 @@ func.func @interface_var(
// spv.target_env
//===----------------------------------------------------------------------===//
func.func @target_env_wrong_limits() attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
// expected-error @+1 {{limits must be a dictionary attribute containing two 32-bit integer attributes 'max_compute_workgroup_invocations' and 'max_compute_workgroup_size'}}
{max_compute_workgroup_invocations = 128 : i64, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
} { return }
// -----
func.func @target_env() attributes {
// CHECK: spv.target_env = #spv.target_env<
// CHECK-SAME: #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
// CHECK-SAME: {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
// CHECK-SAME: #spv.resource_limits<max_compute_workgroup_size = [128, 64, 64]>>
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
{
max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>
}>
#spv.resource_limits<
max_compute_workgroup_size = [128, 64, 64]
>>
} { return }
// -----
@ -131,8 +122,8 @@ func.func @target_env_vendor_id() attributes {
// CHECK: spv.target_env = #spv.target_env<
// CHECK-SAME: #spv.vce<v1.0, [], []>,
// CHECK-SAME: NVIDIA,
// CHECK-SAME: {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, NVIDIA, {}>
// CHECK-SAME: #spv.resource_limits<>>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, NVIDIA, #spv.resource_limits<>>
} { return }
// -----
@ -141,8 +132,8 @@ func.func @target_env_vendor_id_device_type() attributes {
// CHECK: spv.target_env = #spv.target_env<
// CHECK-SAME: #spv.vce<v1.0, [], []>,
// CHECK-SAME: AMD:DiscreteGPU,
// CHECK-SAME: {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, AMD:DiscreteGPU, {}>
// CHECK-SAME: #spv.resource_limits<>>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, AMD:DiscreteGPU, #spv.resource_limits<>>
} { return }
// -----
@ -151,20 +142,17 @@ func.func @target_env_vendor_id_device_type_device_id() attributes {
// CHECK: spv.target_env = #spv.target_env<
// CHECK-SAME: #spv.vce<v1.0, [], []>,
// CHECK-SAME: Qualcomm:IntegratedGPU:100925441,
// CHECK-SAME: {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, Qualcomm:IntegratedGPU:0x6040001, {}>
// CHECK-SAME: #spv.resource_limits<>>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, Qualcomm:IntegratedGPU:0x6040001, #spv.resource_limits<>>
} { return }
// -----
func.func @target_env_extra_fields() attributes {
// expected-error @+6 {{expected '>'}}
// expected-error @+3 {{expected '>'}}
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
{
max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>
},
#spv.resource_limits<>,
more_stuff
>
} { return }
@ -174,37 +162,38 @@ func.func @target_env_extra_fields() attributes {
func.func @target_env_cooperative_matrix() attributes{
// CHECK: spv.target_env = #spv.target_env<
// CHECK-SAME: SPV_NV_cooperative_matrix
// CHECK-SAME: cooperative_matrix_properties_nv = [
// CHECK-SAME: {a_type = i8, b_type = i8, c_type = i32,
// CHECK-SAME: k_size = 32 : i32, m_size = 8 : i32, n_size = 8 : i32
// CHECK-SAME: result_type = i32, scope = 3 : i32}
// CHECK-SAME: {a_type = f16, b_type = f16, c_type = f16,
// CHECK-SAME: k_size = 16 : i32, m_size = 8 : i32, n_size = 8 : i32
// CHECK-SAME: result_type = f16, scope = 3 : i32}
// CHECK-SAME: #spv.coop_matrix_props<
// CHECK-SAME: m_size = 8, n_size = 8, k_size = 32,
// CHECK-SAME: a_type = i8, b_type = i8, c_type = i32,
// CHECK-SAME: result_type = i32, scope = 3 : i32>
// CHECK-SAME: #spv.coop_matrix_props<
// CHECK-SAME: m_size = 8, n_size = 8, k_size = 16,
// CHECK-SAME: a_type = f16, b_type = f16, c_type = f16,
// CHECK-SAME: result_type = f16, scope = 3 : i32>
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class,
SPV_NV_cooperative_matrix]>,
{
cooperative_matrix_properties_nv = [{
m_size = 8: i32,
n_size = 8: i32,
k_size = 32: i32,
#spv.resource_limits<
cooperative_matrix_properties_nv = [#spv.coop_matrix_props<
m_size = 8,
n_size = 8,
k_size = 32,
a_type = i8,
b_type = i8,
c_type = i32,
result_type = i32,
scope = 3: i32
}, {
m_size = 8: i32,
n_size = 8: i32,
k_size = 16: i32,
scope = 3 : i32
>, #spv.coop_matrix_props<
m_size = 8,
n_size = 8,
k_size = 16,
a_type = f16,
b_type = f16,
c_type = f16,
result_type = f16,
scope = 3: i32
}]
}>
scope = 3 : i32
>]
>>
} { return }
// -----

32
mlir/test/Dialect/SPIRV/IR/target-env.mlir
View File

@ -35,7 +35,7 @@
// CHECK-LABEL: @cmp_exchange_weak_suitable_version_capabilities
func.func @cmp_exchange_weak_suitable_version_capabilities(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.1, [Kernel, AtomicStorage], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.1, [Kernel, AtomicStorage], []>, #spv.resource_limits<>>
} {
// CHECK: spv.AtomicCompareExchangeWeak "Workgroup" "AcquireRelease|AtomicCounterMemory" "Acquire"
%0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32)
@ -44,7 +44,7 @@ func.func @cmp_exchange_weak_suitable_version_capabilities(%ptr: !spv.ptr<i32, W
// CHECK-LABEL: @cmp_exchange_weak_unsupported_version
func.func @cmp_exchange_weak_unsupported_version(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.4, [Kernel, AtomicStorage], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.4, [Kernel, AtomicStorage], []>, #spv.resource_limits<>>
} {
// CHECK: test.convert_to_atomic_compare_exchange_weak_op
%0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32)
@ -57,7 +57,7 @@ func.func @cmp_exchange_weak_unsupported_version(%ptr: !spv.ptr<i32, Workgroup>,
// CHECK-LABEL: @group_non_uniform_ballot_suitable_version
func.func @group_non_uniform_ballot_suitable_version(%predicate: i1) -> vector<4xi32> attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.4, [GroupNonUniformBallot], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.4, [GroupNonUniformBallot], []>, #spv.resource_limits<>>
} {
// CHECK: spv.GroupNonUniformBallot Workgroup
%0 = "test.convert_to_group_non_uniform_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
@ -66,7 +66,7 @@ func.func @group_non_uniform_ballot_suitable_version(%predicate: i1) -> vector<4
// CHECK-LABEL: @group_non_uniform_ballot_unsupported_version
func.func @group_non_uniform_ballot_unsupported_version(%predicate: i1) -> vector<4xi32> attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.1, [GroupNonUniformBallot], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.1, [GroupNonUniformBallot], []>, #spv.resource_limits<>>
} {
// CHECK: test.convert_to_group_non_uniform_ballot_op
%0 = "test.convert_to_group_non_uniform_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
@ -79,7 +79,7 @@ func.func @group_non_uniform_ballot_unsupported_version(%predicate: i1) -> vecto
// CHECK-LABEL: @cmp_exchange_weak_missing_capability_kernel
func.func @cmp_exchange_weak_missing_capability_kernel(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.3, [AtomicStorage], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.3, [AtomicStorage], []>, #spv.resource_limits<>>
} {
// CHECK: test.convert_to_atomic_compare_exchange_weak_op
%0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32)
@ -88,7 +88,7 @@ func.func @cmp_exchange_weak_missing_capability_kernel(%ptr: !spv.ptr<i32, Workg
// CHECK-LABEL: @cmp_exchange_weak_missing_capability_atomic_storage
func.func @cmp_exchange_weak_missing_capability_atomic_storage(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.3, [Kernel], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.3, [Kernel], []>, #spv.resource_limits<>>
} {
// CHECK: test.convert_to_atomic_compare_exchange_weak_op
%0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32)
@ -97,7 +97,7 @@ func.func @cmp_exchange_weak_missing_capability_atomic_storage(%ptr: !spv.ptr<i3
// CHECK-LABEL: @subgroup_ballot_missing_capability
func.func @subgroup_ballot_missing_capability(%predicate: i1) -> vector<4xi32> attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.4, [], [SPV_KHR_shader_ballot]>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.4, [], [SPV_KHR_shader_ballot]>, #spv.resource_limits<>>
} {
// CHECK: test.convert_to_subgroup_ballot_op
%0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
@ -106,7 +106,7 @@ func.func @subgroup_ballot_missing_capability(%predicate: i1) -> vector<4xi32> a
// CHECK-LABEL: @bit_reverse_directly_implied_capability
func.func @bit_reverse_directly_implied_capability(%operand: i32) -> i32 attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Geometry], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Geometry], []>, #spv.resource_limits<>>
} {
// CHECK: spv.BitReverse
%0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32)
@ -115,7 +115,7 @@ func.func @bit_reverse_directly_implied_capability(%operand: i32) -> i32 attribu
// CHECK-LABEL: @bit_reverse_recursively_implied_capability
func.func @bit_reverse_recursively_implied_capability(%operand: i32) -> i32 attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [GeometryPointSize], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [GeometryPointSize], []>, #spv.resource_limits<>>
} {
// CHECK: spv.BitReverse
%0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32)
@ -128,7 +128,7 @@ func.func @bit_reverse_recursively_implied_capability(%operand: i32) -> i32 attr
// CHECK-LABEL: @subgroup_ballot_suitable_extension
func.func @subgroup_ballot_suitable_extension(%predicate: i1) -> vector<4xi32> attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], [SPV_KHR_shader_ballot]>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], [SPV_KHR_shader_ballot]>, #spv.resource_limits<>>
} {
// CHECK: spv.SubgroupBallotKHR
%0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
@ -137,7 +137,7 @@ func.func @subgroup_ballot_suitable_extension(%predicate: i1) -> vector<4xi32> a
// CHECK-LABEL: @subgroup_ballot_missing_extension
func.func @subgroup_ballot_missing_extension(%predicate: i1) -> vector<4xi32> attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], []>, #spv.resource_limits<>>
} {
// CHECK: test.convert_to_subgroup_ballot_op
%0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
@ -146,7 +146,7 @@ func.func @subgroup_ballot_missing_extension(%predicate: i1) -> vector<4xi32> at
// CHECK-LABEL: @module_suitable_extension1
func.func @module_suitable_extension1() attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_EXT_physical_storage_buffer]>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_EXT_physical_storage_buffer]>, #spv.resource_limits<>>
} {
// CHECK: spv.module PhysicalStorageBuffer64 Vulkan
"test.convert_to_module_op"() : () ->()
@ -155,7 +155,7 @@ func.func @module_suitable_extension1() attributes {
// CHECK-LABEL: @module_suitable_extension2
func.func @module_suitable_extension2() attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_KHR_physical_storage_buffer]>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_KHR_physical_storage_buffer]>, #spv.resource_limits<>>
} {
// CHECK: spv.module PhysicalStorageBuffer64 Vulkan
"test.convert_to_module_op"() : () -> ()
@ -164,7 +164,7 @@ func.func @module_suitable_extension2() attributes {
// CHECK-LABEL: @module_missing_extension_mm
func.func @module_missing_extension_mm() attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_physical_storage_buffer]>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_physical_storage_buffer]>, #spv.resource_limits<>>
} {
// CHECK: test.convert_to_module_op
"test.convert_to_module_op"() : () -> ()
@ -173,7 +173,7 @@ func.func @module_missing_extension_mm() attributes {
// CHECK-LABEL: @module_missing_extension_am
func.func @module_missing_extension_am() attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model]>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model]>, #spv.resource_limits<>>
} {
// CHECK: test.convert_to_module_op
"test.convert_to_module_op"() : () -> ()
@ -183,7 +183,7 @@ func.func @module_missing_extension_am() attributes {
// CHECK-LABEL: @module_implied_extension
func.func @module_implied_extension() attributes {
// Version 1.5 implies SPV_KHR_vulkan_memory_model and SPV_KHR_physical_storage_buffer.
spv.target_env = #spv.target_env<#spv.vce<v1.5, [VulkanMemoryModel, PhysicalStorageBufferAddresses], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.5, [VulkanMemoryModel, PhysicalStorageBufferAddresses], []>, #spv.resource_limits<>>
} {
// CHECK: spv.module PhysicalStorageBuffer64 Vulkan
"test.convert_to_module_op"() : () -> ()

4
mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir
View File

@ -259,14 +259,14 @@ spv.module Logical GLSL450 {
spv.func @kernel(
%arg0: f32,
%arg1: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, CrossWorkgroup>) "None"
attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
spv.Return
}
spv.func @kernel_different_attr(
%arg0: f32,
%arg1: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, CrossWorkgroup>) "None"
attributes {spv.entry_point_abi = {local_size = dense<[64, 1, 1]> : vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[64, 1, 1]> : vector<3xi32>>} {
spv.Return
}
}

4
mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir
View File

@ -1,7 +1,7 @@
// RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
module attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel, Addresses], []>, {}>
spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel, Addresses], []>, #spv.resource_limits<>>
} {
spv.module Physical64 OpenCL {
// CHECK-LABEL: spv.module
@ -11,7 +11,7 @@ module attributes {
spv.func @kernel(
%arg0: f32,
%arg1: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, CrossWorkgroup>) "None"
attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
spv.Return
}
}

4
mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir
View File

@ -2,7 +2,7 @@
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.module
@ -15,7 +15,7 @@ spv.module Logical GLSL450 {
{spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>},
%arg1: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, StorageBuffer>
{spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}) "None"
attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
// CHECK: [[ARG1:%.*]] = spv.mlir.addressof [[VAR1]]
// CHECK: [[ADDRESSARG0:%.*]] = spv.mlir.addressof [[VAR0]]
// CHECK: [[CONST0:%.*]] = spv.Constant 0 : i32

4
mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir
View File

@ -2,7 +2,7 @@
module attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
// CHECK-LABEL: spv.module
@ -38,7 +38,7 @@ spv.module Logical GLSL450 {
{spv.interface_var_abi = #spv.interface_var_abi<(0, 5), StorageBuffer>},
%arg6: i32
{spv.interface_var_abi = #spv.interface_var_abi<(0, 6), StorageBuffer>}) "None"
attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
// CHECK: [[ADDRESSARG6:%.*]] = spv.mlir.addressof [[VAR6]]
// CHECK: [[CONST6:%.*]] = spv.Constant 0 : i32
// CHECK: [[ARG6PTR:%.*]] = spv.AccessChain [[ADDRESSARG6]]{{\[}}[[CONST6]]

28
mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir
View File

@ -10,7 +10,7 @@
// CHECK: requires #spv.vce<v1.0, [Shader], []>
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.5, [Shader], []>, {}>
#spv.vce<v1.5, [Shader], []>, #spv.resource_limits<>>
} {
spv.func @iadd(%val : i32) -> i32 "None" {
%0 = spv.IAdd %val, %val: i32
@ -24,7 +24,7 @@ spv.module Logical GLSL450 attributes {
// CHECK: requires #spv.vce<v1.3, [GroupNonUniformBallot, Shader], []>
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.5, [Shader, GroupNonUniformBallot], []>, {}>
#spv.vce<v1.5, [Shader, GroupNonUniformBallot], []>, #spv.resource_limits<>>
} {
spv.func @group_non_uniform_ballot(%predicate : i1) -> vector<4xi32> "None" {
%0 = spv.GroupNonUniformBallot Workgroup %predicate : vector<4xi32>
@ -41,7 +41,7 @@ spv.module Logical GLSL450 attributes {
// CHECK: requires #spv.vce<v1.0, [Shader], []>
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader, Float16, Float64, Int16, Int64, VariablePointers], []>, {}>
#spv.vce<v1.0, [Shader, Float16, Float64, Int16, Int64, VariablePointers], []>, #spv.resource_limits<>>
} {
spv.func @iadd(%val : i32) -> i32 "None" {
%0 = spv.IAdd %val, %val: i32
@ -54,7 +54,7 @@ spv.module Logical GLSL450 attributes {
// CHECK: spv.module PhysicalStorageBuffer64 GLSL450 requires #spv.vce<v1.0, [PhysicalStorageBufferAddresses, Shader], [SPV_EXT_physical_storage_buffer]>
spv.module PhysicalStorageBuffer64 GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader, PhysicalStorageBufferAddresses], [SPV_EXT_physical_storage_buffer]>, {}>
#spv.vce<v1.0, [Shader, PhysicalStorageBufferAddresses], [SPV_EXT_physical_storage_buffer]>, #spv.resource_limits<>>
} {
spv.func @physical_ptr(%val : !spv.ptr<f32, PhysicalStorageBuffer>) "None" {
spv.Return
@ -67,7 +67,7 @@ spv.module PhysicalStorageBuffer64 GLSL450 attributes {
// CHECK: requires #spv.vce<v1.0, [Shader], []>
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [AtomicStorage], []>, {}>
#spv.vce<v1.0, [AtomicStorage], []>, #spv.resource_limits<>>
} {
spv.func @iadd(%val : i32) -> i32 "None" {
%0 = spv.IAdd %val, %val: i32
@ -88,7 +88,7 @@ spv.module Logical GLSL450 attributes {
// CHECK: requires #spv.vce<v1.3, [GroupNonUniformArithmetic, Shader], []>
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
#spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #spv.resource_limits<>>
} {
spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" {
%0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32
@ -99,7 +99,7 @@ spv.module Logical GLSL450 attributes {
// CHECK: requires #spv.vce<v1.3, [GroupNonUniformClustered, GroupNonUniformBallot, Shader], []>
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, GroupNonUniformClustered, GroupNonUniformBallot], []>, {}>
#spv.vce<v1.3, [Shader, GroupNonUniformClustered, GroupNonUniformBallot], []>, #spv.resource_limits<>>
} {
spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" {
%0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32
@ -113,7 +113,7 @@ spv.module Logical GLSL450 attributes {
// CHECK: requires #spv.vce<v1.0, [Int8, Shader], []>
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, Int8], []>, {}>
#spv.vce<v1.3, [Shader, Int8], []>, #spv.resource_limits<>>
} {
spv.func @iadd_function(%val : i8) -> i8 "None" {
%0 = spv.IAdd %val, %val : i8
@ -125,7 +125,7 @@ spv.module Logical GLSL450 attributes {
// CHECK: requires #spv.vce<v1.0, [Float16, Shader], []>
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, Float16], []>, {}>
#spv.vce<v1.3, [Shader, Float16], []>, #spv.resource_limits<>>
} {
spv.func @fadd_function(%val : f16) -> f16 "None" {
%0 = spv.FAdd %val, %val : f16
@ -137,7 +137,7 @@ spv.module Logical GLSL450 attributes {
// CHECK: requires #spv.vce<v1.0, [Vector16, Shader], []>
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, Vector16], []>, {}>
#spv.vce<v1.3, [Shader, Vector16], []>, #spv.resource_limits<>>
} {
spv.func @iadd_v16_function(%val : vector<16xi32>) -> vector<16xi32> "None" {
%0 = spv.IAdd %val, %val : vector<16xi32>
@ -156,7 +156,7 @@ spv.module Logical GLSL450 attributes {
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader, SubgroupBallotKHR],
[SPV_KHR_shader_ballot, SPV_KHR_shader_clock, SPV_KHR_variable_pointers]>, {}>
[SPV_KHR_shader_ballot, SPV_KHR_shader_clock, SPV_KHR_variable_pointers]>, #spv.resource_limits<>>
} {
spv.func @subgroup_ballot(%predicate : i1) -> vector<4xi32> "None" {
%0 = spv.SubgroupBallotKHR %predicate: vector<4xi32>
@ -171,7 +171,7 @@ spv.module Logical GLSL450 attributes {
// CHECK: requires #spv.vce<v1.0, [VulkanMemoryModel], [SPV_KHR_vulkan_memory_model]>
spv.module Logical Vulkan attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.5, [Shader, VulkanMemoryModel], []>, {}>
#spv.vce<v1.5, [Shader, VulkanMemoryModel], []>, #spv.resource_limits<>>
} {
spv.func @iadd(%val : i32) -> i32 "None" {
%0 = spv.IAdd %val, %val: i32
@ -186,7 +186,7 @@ spv.module Logical Vulkan attributes {
// CHECK: requires #spv.vce<v1.0, [StorageBuffer16BitAccess, Shader, Int16], [SPV_KHR_16bit_storage, SPV_KHR_storage_buffer_storage_class]>
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, StorageBuffer16BitAccess, Int16], []>, {}>
#spv.vce<v1.3, [Shader, StorageBuffer16BitAccess, Int16], []>, #spv.resource_limits<>>
} {
spv.func @iadd_storage_buffer(%ptr : !spv.ptr<i16, StorageBuffer>) -> i16 "None" {
%0 = spv.Load "StorageBuffer" %ptr : i16
@ -202,7 +202,7 @@ spv.module Logical GLSL450 attributes {
spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.5, [Shader, UniformAndStorageBuffer8BitAccess, StorageBuffer16BitAccess, StorageUniform16, Int16, Int64, ImageBuffer, StorageImageExtendedFormats], []>,
{}>
#spv.resource_limits<>>
} {
spv.GlobalVariable @data : !spv.ptr<!spv.struct<(i8 [0], f16 [2], i64 [4])>, Uniform>
spv.GlobalVariable @img : !spv.ptr<!spv.image<f32, Buffer, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Rg32f>, UniformConstant>

3
mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp
View File

@ -31,6 +31,9 @@ struct TestSpirvEntryPointABIPass
"within the "
"module, intended for testing only";
}
void getDependentDialects(DialectRegistry &registry) const override {
registry.insert<spirv::SPIRVDialect>();
}
TestSpirvEntryPointABIPass() = default;
TestSpirvEntryPointABIPass(const TestSpirvEntryPointABIPass &) {}
void runOnOperation() override;

7
mlir/test/mlir-spirv-cpu-runner/double.mlir
View File

@ -5,12 +5,13 @@ module attributes {
gpu.container_module,
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>,
{max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
#spv.resource_limits<
max_compute_workgroup_invocations = 128,
max_compute_workgroup_size = [128, 128, 64]>>
} {
gpu.module @kernels {
gpu.func @double(%arg0 : memref<6xi32>, %arg1 : memref<6xi32>)
kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32>}} {
kernel attributes { spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
%factor = arith.constant 2 : i32
%i0 = arith.constant 0 : index

7
mlir/test/mlir-spirv-cpu-runner/simple_add.mlir
View File

@ -5,12 +5,13 @@ module attributes {
gpu.container_module,
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_8bit_storage]>,
{max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
#spv.resource_limits<
max_compute_workgroup_invocations = 128,
max_compute_workgroup_size = [128, 128, 64]>>
} {
gpu.module @kernels {
gpu.func @sum(%arg0 : memref<3xf32>, %arg1 : memref<3x3xf32>, %arg2 : memref<3x3x3xf32>)
kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32>}} {
kernel attributes { spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
%i0 = arith.constant 0 : index
%i1 = arith.constant 1 : index
%i2 = arith.constant 2 : index

4
mlir/test/mlir-vulkan-runner/addf.mlir
View File

@ -4,11 +4,11 @@
module attributes {
gpu.container_module,
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
gpu.module @kernels {
gpu.func @kernel_add(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32> }} {
kernel attributes { spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
%0 = gpu.block_id x
%1 = memref.load %arg0[%0] : memref<8xf32>
%2 = memref.load %arg1[%0] : memref<8xf32>

4
mlir/test/mlir-vulkan-runner/addi.mlir
View File

@ -4,11 +4,11 @@
module attributes {
gpu.container_module,
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
gpu.module @kernels {
gpu.func @kernel_addi(%arg0 : memref<8xi32>, %arg1 : memref<8x8xi32>, %arg2 : memref<8x8x8xi32>)
kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32>}} {
kernel attributes { spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
%x = gpu.block_id x
%y = gpu.block_id y
%z = gpu.block_id z

4
mlir/test/mlir-vulkan-runner/addi8.mlir
View File

@ -4,11 +4,11 @@
module attributes {
gpu.container_module,
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_8bit_storage]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_8bit_storage]>, #spv.resource_limits<>>
} {
gpu.module @kernels {
gpu.func @kernel_addi(%arg0 : memref<8xi8>, %arg1 : memref<8x8xi8>, %arg2 : memref<8x8x8xi32>)
kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32>}} {
kernel attributes { spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
%x = gpu.block_id x
%y = gpu.block_id y
%z = gpu.block_id z

4
mlir/test/mlir-vulkan-runner/mulf.mlir
View File

@ -4,11 +4,11 @@
module attributes {
gpu.container_module,
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
} {
gpu.module @kernels {
gpu.func @kernel_mul(%arg0 : memref<4x4xf32>, %arg1 : memref<4x4xf32>, %arg2 : memref<4x4xf32>)
kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32> }} {
kernel attributes { spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
%x = gpu.block_id x
%y = gpu.block_id y
%1 = memref.load %arg0[%x, %y] : memref<4x4xf32>

5
mlir/test/mlir-vulkan-runner/subf.mlir
View File

@ -4,11 +4,12 @@
module attributes {
gpu.container_module,
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
#spv.resource_limits<>>
} {
gpu.module @kernels {
gpu.func @kernel_sub(%arg0 : memref<8x4x4xf32>, %arg1 : memref<4x4xf32>, %arg2 : memref<8x4x4xf32>)
kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32> }} {
kernel attributes { spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
%x = gpu.block_id x
%y = gpu.block_id y
%z = gpu.block_id z

4
mlir/test/mlir-vulkan-runner/time.mlir
View File

@ -7,11 +7,11 @@
module attributes {
gpu.container_module,
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits>
} {
gpu.module @kernels {
gpu.func @kernel_add(%arg0 : memref<16384xf32>, %arg1 : memref<16384xf32>, %arg2 : memref<16384xf32>)
kernel attributes { spv.entry_point_abi = {local_size = dense<[128, 1, 1]>: vector<3xi32> }} {
kernel attributes { spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[128, 1, 1]>: vector<3xi32>>} {
%bid = gpu.block_id x
%tid = gpu.thread_id x
%cst = arith.constant 128 : index

4
mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp
View File

@ -250,9 +250,11 @@ void DefFormat::genParser(MethodBody &os) {
FmtContext ctx;
ctx.addSubst("_parser", "odsParser");
ctx.addSubst("_ctx", "odsParser.getContext()");
ctx.withBuilder("odsBuilder");
if (isa<AttrDef>(def))
ctx.addSubst("_type", "odsType");
os.indent();
os << "::mlir::Builder odsBuilder(odsParser.getContext());\n";
// Declare variables to store all of the parameters. Allocated parameters
// such as `ArrayRef` and `StringRef` must provide a `storageType`. Store
@ -668,7 +670,9 @@ void DefFormat::genPrinter(MethodBody &os) {
FmtContext ctx;
ctx.addSubst("_printer", "odsPrinter");
ctx.addSubst("_ctx", "getContext()");
ctx.withBuilder("odsBuilder");
os.indent();
os << "::mlir::Builder odsBuilder(getContext());\n";
// Generate printers.
shouldEmitSpace = true;

37
utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
View File

@ -4471,6 +4471,24 @@ gentbl_cc_library(
deps = [":SPIRVOpsTdFiles"],
)
gentbl_cc_library(
name = "SPIRVAttributesIncGen",
strip_include_prefix = "include",
tbl_outs = [
(
["-gen-attrdef-decls"],
"include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.h.inc",
),
(
["-gen-attrdef-defs"],
"include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.cpp.inc",
),
],
tblgen = ":mlir-tblgen",
td_file = "include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.td",
deps = [":SPIRVOpsTdFiles"],
)
gentbl_cc_library(
name = "SPIRVCanonicalizationIncGen",
strip_include_prefix = "lib/Dialect/SPIRV/IR",
@ -4507,23 +4525,6 @@ gentbl_cc_library(
deps = [":SPIRVOpsTdFiles"],
)
gentbl_cc_library(
name = "SPIRVTargetAndABIStructGen",
tbl_outs = [
(
["-gen-struct-attr-decls"],
"include/mlir/Dialect/SPIRV/IR/TargetAndABI.h.inc",
),
(
["-gen-struct-attr-defs"],
"include/mlir/Dialect/SPIRV/IR/TargetAndABI.cpp.inc",
),
],
tblgen = ":mlir-tblgen",
td_file = "include/mlir/Dialect/SPIRV/IR/TargetAndABI.td",
deps = [":SPIRVOpsTdFiles"],
)
gentbl_cc_library(
name = "SPIRVAttrUtilsGen",
strip_include_prefix = "include",
@ -4569,11 +4570,11 @@ cc_library(
":InferTypeOpInterface",
":Parser",
":SPIRVAttrUtilsGen",
":SPIRVAttributesIncGen",
":SPIRVAvailabilityIncGen",
":SPIRVCanonicalizationIncGen",
":SPIRVOpsIncGen",
":SPIRVSerializationGen",
":SPIRVTargetAndABIStructGen",
":SideEffectInterfaces",
":Support",
":Transforms",