[Mlir-commits] [mlir] e115a40 - [mlir][spirv] Use separate attribute for (version, capabilities, extensions)
Lei Zhang
llvmlistbot at llvm.org
Thu Mar 12 16:39:46 PDT 2020
Author: Lei Zhang
Date: 2020-03-12T19:37:45-04:00
New Revision: e115a40f5020b188c01ecfd435b02e7e46393b0d
URL: https://github.com/llvm/llvm-project/commit/e115a40f5020b188c01ecfd435b02e7e46393b0d
DIFF: https://github.com/llvm/llvm-project/commit/e115a40f5020b188c01ecfd435b02e7e46393b0d.diff
LOG: [mlir][spirv] Use separate attribute for (version, capabilities, extensions)
We also need the (version, capabilities, extensions) triple on the
spv.module op. Thus far we have been using separate 'extensions'
and 'capabilities' attributes there and 'version' is missing. Creating
a separate attribute for the trip allows us to reuse the assembly
form and verification.
Differential Revision: https://reviews.llvm.org/D75868
Added:
Modified:
mlir/docs/Dialects/SPIR-V.md
mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td
mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h
mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp
mlir/lib/Dialect/SPIRV/TargetAndABI.cpp
mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir
mlir/test/Dialect/SPIRV/availability.mlir
mlir/test/Dialect/SPIRV/target-and-abi.mlir
mlir/test/Dialect/SPIRV/target-env.mlir
Removed:
################################################################################
diff --git a/mlir/docs/Dialects/SPIR-V.md b/mlir/docs/Dialects/SPIR-V.md
index d81d1038ca43..4380c9f006d1 100644
--- a/mlir/docs/Dialects/SPIR-V.md
+++ b/mlir/docs/Dialects/SPIR-V.md
@@ -742,11 +742,11 @@ instructions.
SPIR-V compilation should also take into consideration of the execution
environment, so we generate SPIR-V modules valid for the target environment.
-This is conveyed by the `spv.target_env` attribute. It should be of
-`#spv.target_env` attribute kind, which is defined as:
+This is conveyed by the `spv.target_env` (`spirv::TargetEnvAttr`) attribute. It
+should be of `#spv.target_env` attribute kind, which is defined as:
```
-spirv-version ::= `V_1_0` | `V_1_1` | ...
+spirv-version ::= `v1.0` | `v1.1` | ...
spirv-extension ::= `SPV_KHR_16bit_storage` | `SPV_EXT_physical_storage_buffer` | ...
spirv-capability ::= `Shader` | `Kernel` | `GroupNonUniform` | ...
@@ -758,18 +758,22 @@ spirv-capability-elements ::= spirv-capability (`,` spirv-capability)*
spirv-resource-limits ::= dictionary-attribute
+spirv-vce-attribute ::= `#` `spv.vce` `<`
+ spirv-version `,`
+ spirv-capability-list `,`
+ spirv-extensions-list `>`
+
spirv-target-env-attribute ::= `#` `spv.target_env` `<`
- spirv-version `,`
- spirv-extensions-list `,`
- spirv-capability-list `,`
+ spirv-vce-attribute,
spirv-resource-limits `>`
```
The attribute has a few fields:
-* The target SPIR-V version.
-* A list of SPIR-V extensions for the target.
-* A list of SPIR-V capabilities for the target.
+* A `#spv.vce` (`spirv::VerCapExtAttr`) attribute:
+ * The target SPIR-V version.
+ * A list of SPIR-V extensions for the target.
+ * A list of SPIR-V capabilities for the target.
* A dictionary of target resource limits (see the
[Vulkan spec][VulkanResourceLimits] for explanation):
* `max_compute_workgroup_invocations`
@@ -780,7 +784,7 @@ For example,
```
module attributes {
spv.target_env = #spv.target_env<
- V_1_3, [SPV_KHR_8bit_storage], [Shader, GroupNonUniform]
+ #spv.vce<v1.3, [Shader, GroupNonUniform], [SPV_KHR_8bit_storage]>,
{
max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>
diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td
index 4eefc6189903..1af6ddef4ea0 100644
--- a/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td
+++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td
@@ -96,12 +96,12 @@ class SPV_StrEnumAttr<string name, string description,
// SPIR-V availability definitions
//===----------------------------------------------------------------------===//
-def SPV_V_1_0 : I32EnumAttrCase<"V_1_0", 0>;
-def SPV_V_1_1 : I32EnumAttrCase<"V_1_1", 1>;
-def SPV_V_1_2 : I32EnumAttrCase<"V_1_2", 2>;
-def SPV_V_1_3 : I32EnumAttrCase<"V_1_3", 3>;
-def SPV_V_1_4 : I32EnumAttrCase<"V_1_4", 4>;
-def SPV_V_1_5 : I32EnumAttrCase<"V_1_5", 5>;
+def SPV_V_1_0 : I32EnumAttrCase<"V_1_0", 0, "v1.0">;
+def SPV_V_1_1 : I32EnumAttrCase<"V_1_1", 1, "v1.1">;
+def SPV_V_1_2 : I32EnumAttrCase<"V_1_2", 2, "v1.2">;
+def SPV_V_1_3 : I32EnumAttrCase<"V_1_3", 3, "v1.3">;
+def SPV_V_1_4 : I32EnumAttrCase<"V_1_4", 4, "v1.4">;
+def SPV_V_1_5 : I32EnumAttrCase<"V_1_5", 5, "v1.5">;
def SPV_VersionAttr : SPV_I32EnumAttr<"Version", "valid SPIR-V version", [
SPV_V_1_0, SPV_V_1_1, SPV_V_1_2, SPV_V_1_3, SPV_V_1_4, SPV_V_1_5]>;
diff --git a/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h b/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h
index 01b775846ee3..1d3964a67fe5 100644
--- a/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h
+++ b/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h
@@ -32,36 +32,37 @@ enum class Version : uint32_t;
namespace detail {
struct TargetEnvAttributeStorage;
+struct VerCapExtAttributeStorage;
} // namespace detail
/// SPIR-V dialect-specific attribute kinds.
// TODO(antiagainst): move to a more suitable place if we have more attributes.
namespace AttrKind {
enum Kind {
- TargetEnv = Attribute::FIRST_SPIRV_ATTR,
+ TargetEnv = Attribute::FIRST_SPIRV_ATTR, /// Target environment
+ VerCapExt, /// (version, extension, capability) triple
};
} // namespace AttrKind
-/// An attribute that specifies the target version, allowed extensions and
-/// capabilities, and resource limits. These information describles a SPIR-V
-/// target environment.
-class TargetEnvAttr
- : public Attribute::AttrBase<TargetEnvAttr, Attribute,
- detail::TargetEnvAttributeStorage> {
+/// An attribute that specifies the SPIR-V (version, capabilities, extensions)
+/// triple.
+class VerCapExtAttr
+ : public Attribute::AttrBase<VerCapExtAttr, Attribute,
+ detail::VerCapExtAttributeStorage> {
public:
using Base::Base;
- /// Gets a TargetEnvAttr instance.
- static TargetEnvAttr get(Version version, ArrayRef<Extension> extensions,
- ArrayRef<Capability> capabilities,
- DictionaryAttr limits);
- static TargetEnvAttr get(IntegerAttr version, ArrayAttr extensions,
- ArrayAttr capabilities, DictionaryAttr limits);
+ /// Gets a VerCapExtAttr instance.
+ static VerCapExtAttr get(Version version, ArrayRef<Capability> capabilities,
+ ArrayRef<Extension> extensions,
+ MLIRContext *context);
+ static VerCapExtAttr get(IntegerAttr version, ArrayAttr capabilities,
+ ArrayAttr extensions);
/// Returns the attribute kind's name (without the 'spv.' prefix).
static StringRef getKindName();
- /// Returns the target version.
+ /// Returns the version.
Version getVersion();
struct ext_iterator final
@@ -71,9 +72,9 @@ class TargetEnvAttr
};
using ext_range = llvm::iterator_range<ext_iterator>;
- /// Returns the target extensions.
+ /// Returns the extensions.
ext_range getExtensions();
- /// Returns the target extensions as a string array attribute.
+ /// Returns the extensions as a string array attribute.
ArrayAttr getExtensionsAttr();
struct cap_iterator final
@@ -83,8 +84,47 @@ class TargetEnvAttr
};
using cap_range = llvm::iterator_range<cap_iterator>;
- /// Returns the target capabilities.
+ /// Returns the capabilities.
cap_range getCapabilities();
+ /// Returns the capabilities as an integer array attribute.
+ ArrayAttr getCapabilitiesAttr();
+
+ static bool kindof(unsigned kind) { return kind == AttrKind::VerCapExt; }
+
+ static LogicalResult verifyConstructionInvariants(Location loc,
+ IntegerAttr version,
+ ArrayAttr capabilities,
+ ArrayAttr extensions);
+};
+
+/// An attribute that specifies the target version, allowed extensions and
+/// capabilities, and resource limits. These information describles a SPIR-V
+/// target environment.
+class TargetEnvAttr
+ : public Attribute::AttrBase<TargetEnvAttr, Attribute,
+ detail::TargetEnvAttributeStorage> {
+public:
+ using Base::Base;
+
+ /// Gets a TargetEnvAttr instance.
+ static TargetEnvAttr get(VerCapExtAttr triple, DictionaryAttr limits);
+
+ /// Returns the attribute kind's name (without the 'spv.' prefix).
+ static StringRef getKindName();
+
+ /// Returns the (version, capabilities, extensions) triple attribute.
+ VerCapExtAttr getTripleAttr();
+
+ /// Returns the target version.
+ Version getVersion();
+
+ /// Returns the target extensions.
+ VerCapExtAttr::ext_range getExtensions();
+ /// Returns the target extensions as a string array attribute.
+ ArrayAttr getExtensionsAttr();
+
+ /// Returns the target capabilities.
+ VerCapExtAttr::cap_range getCapabilities();
/// Returns the target capabilities as an integer array attribute.
ArrayAttr getCapabilitiesAttr();
@@ -94,9 +134,7 @@ class TargetEnvAttr
static bool kindof(unsigned kind) { return kind == AttrKind::TargetEnv; }
static LogicalResult verifyConstructionInvariants(Location loc,
- IntegerAttr version,
- ArrayAttr extensions,
- ArrayAttr capabilities,
+ VerCapExtAttr triple,
DictionaryAttr limits);
};
diff --git a/mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp b/mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp
index 1946bfc37ce3..50ecf9ef7cbd 100644
--- a/mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp
+++ b/mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp
@@ -118,7 +118,7 @@ SPIRVDialect::SPIRVDialect(MLIRContext *context)
: Dialect(getDialectNamespace(), context) {
addTypes<ArrayType, ImageType, PointerType, RuntimeArrayType, StructType>();
- addAttributes<TargetEnvAttr>();
+ addAttributes<TargetEnvAttr, VerCapExtAttr>();
// Add SPIR-V ops.
addOperations<
@@ -662,8 +662,7 @@ static ParseResult parseKeywordList(
return success();
}
-/// Parses a spirv::TargetEnvAttr.
-static Attribute parseTargetAttr(DialectAsmParser &parser) {
+static Attribute parseVerCapExtAttr(DialectAsmParser &parser) {
if (parser.parseLess())
return {};
@@ -685,51 +684,67 @@ static Attribute parseTargetAttr(DialectAsmParser &parser) {
}
}
- ArrayAttr extensionsAttr;
+ ArrayAttr capabilitiesAttr;
{
- SmallVector<Attribute, 1> extensions;
+ SmallVector<Attribute, 4> capabilities;
llvm::SMLoc errorloc;
StringRef errorKeyword;
- auto processExtension = [&](llvm::SMLoc loc, StringRef extension) {
- if (spirv::symbolizeExtension(extension)) {
- extensions.push_back(builder.getStringAttr(extension));
+ auto processCapability = [&](llvm::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 = extension, failure();
+ return errorloc = loc, errorKeyword = capability, failure();
};
- if (parseKeywordList(parser, processExtension) || parser.parseComma()) {
+ if (parseKeywordList(parser, processCapability) || parser.parseComma()) {
if (!errorKeyword.empty())
- parser.emitError(errorloc, "unknown extension: ") << errorKeyword;
+ parser.emitError(errorloc, "unknown capability: ") << errorKeyword;
return {};
}
- extensionsAttr = builder.getArrayAttr(extensions);
+ capabilitiesAttr = builder.getArrayAttr(capabilities);
}
- ArrayAttr capabilitiesAttr;
+ ArrayAttr extensionsAttr;
{
- SmallVector<Attribute, 4> capabilities;
+ SmallVector<Attribute, 1> extensions;
llvm::SMLoc errorloc;
StringRef errorKeyword;
- auto processCapability = [&](llvm::SMLoc loc, StringRef capability) {
- if (auto capSymbol = spirv::symbolizeCapability(capability)) {
- capabilities.push_back(
- builder.getI32IntegerAttr(static_cast<uint32_t>(*capSymbol)));
+ auto processExtension = [&](llvm::SMLoc loc, StringRef extension) {
+ if (spirv::symbolizeExtension(extension)) {
+ extensions.push_back(builder.getStringAttr(extension));
return success();
}
- return errorloc = loc, errorKeyword = capability, failure();
+ return errorloc = loc, errorKeyword = extension, failure();
};
- if (parseKeywordList(parser, processCapability) || parser.parseComma()) {
+ if (parseKeywordList(parser, processExtension)) {
if (!errorKeyword.empty())
- parser.emitError(errorloc, "unknown capability: ") << errorKeyword;
+ parser.emitError(errorloc, "unknown extension: ") << errorKeyword;
return {};
}
- capabilitiesAttr = builder.getArrayAttr(capabilities);
+ 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 {};
+
DictionaryAttr limitsAttr;
{
auto loc = parser.getCurrentLocation();
@@ -749,8 +764,7 @@ static Attribute parseTargetAttr(DialectAsmParser &parser) {
if (parser.parseGreater())
return {};
- return spirv::TargetEnvAttr::get(versionAttr, extensionsAttr,
- capabilitiesAttr, limitsAttr);
+ return spirv::TargetEnvAttr::get(tripleAttr, limitsAttr);
}
Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser,
@@ -767,7 +781,9 @@ Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser,
return {};
if (attrKind == spirv::TargetEnvAttr::getKindName())
- return parseTargetAttr(parser);
+ return parseTargetEnvAttr(parser);
+ if (attrKind == spirv::VerCapExtAttr::getKindName())
+ return parseVerCapExtAttr(parser);
parser.emitError(parser.getNameLoc(), "unknown SPIR-V attriubte kind: ")
<< attrKind;
@@ -778,24 +794,32 @@ Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser,
// Attribute Printing
//===----------------------------------------------------------------------===//
-static void print(spirv::TargetEnvAttr targetEnv, DialectAsmPrinter &printer) {
+static void print(spirv::VerCapExtAttr triple, DialectAsmPrinter &printer) {
auto &os = printer.getStream();
- printer << spirv::TargetEnvAttr::getKindName() << "<"
- << spirv::stringifyVersion(targetEnv.getVersion()) << ", [";
- interleaveComma(targetEnv.getExtensionsAttr(), os, [&](Attribute attr) {
- os << attr.cast<StringAttr>().getValue();
+ printer << spirv::VerCapExtAttr::getKindName() << "<"
+ << spirv::stringifyVersion(triple.getVersion()) << ", [";
+ interleaveComma(triple.getCapabilities(), os, [&](spirv::Capability cap) {
+ os << spirv::stringifyCapability(cap);
});
printer << "], [";
- interleaveComma(targetEnv.getCapabilities(), os, [&](spirv::Capability cap) {
- os << spirv::stringifyCapability(cap);
+ interleaveComma(triple.getExtensionsAttr(), os, [&](Attribute attr) {
+ os << attr.cast<StringAttr>().getValue();
});
- printer << "], " << targetEnv.getResourceLimits() << ">";
+ printer << "]>";
+}
+
+static void print(spirv::TargetEnvAttr targetEnv, DialectAsmPrinter &printer) {
+ printer << spirv::TargetEnvAttr::getKindName() << "<#spv.";
+ print(targetEnv.getTripleAttr(), printer);
+ printer << ", " << targetEnv.getResourceLimits() << ">";
}
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
llvm_unreachable("unhandled SPIR-V attribute kind");
}
@@ -807,7 +831,7 @@ void SPIRVDialect::printAttribute(Attribute attr,
Operation *SPIRVDialect::materializeConstant(OpBuilder &builder,
Attribute value, Type type,
Location loc) {
- if (!ConstantOp::isBuildableWith(type))
+ if (!spirv::ConstantOp::isBuildableWith(type))
return nullptr;
return builder.create<spirv::ConstantOp>(loc, type, value);
@@ -832,12 +856,7 @@ LogicalResult SPIRVDialect::verifyOperationAttribute(Operation *op,
"32-bit integer elements attribute: 'local_size'";
} else if (symbol == spirv::getTargetEnvAttrName()) {
if (!attr.isa<spirv::TargetEnvAttr>())
- return op->emitError("'")
- << symbol
- << "' must be a dictionary attribute containing one 32-bit "
- "integer attribute 'version', one string array attribute "
- "'extensions', one 32-bit integer array attribute "
- "'capabilities', and one dictionary attribute 'limits'";
+ return op->emitError("'") << symbol << "' must be a spirv::TargetEnvAttr";
} else {
return op->emitError("found unsupported '")
<< symbol << "' attribute on operation";
diff --git a/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp b/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp
index f8c5900eb842..3743cf44348c 100644
--- a/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp
+++ b/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp
@@ -14,95 +14,123 @@
using namespace mlir;
+//===----------------------------------------------------------------------===//
+// DictionaryDict derived attributes
+//===----------------------------------------------------------------------===//
+
namespace mlir {
#include "mlir/Dialect/SPIRV/TargetAndABI.cpp.inc"
+//===----------------------------------------------------------------------===//
+// Attribute storage classes
+//===----------------------------------------------------------------------===//
+
namespace spirv {
namespace detail {
+struct VerCapExtAttributeStorage : public AttributeStorage {
+ using KeyTy = std::tuple<Attribute, Attribute, Attribute>;
+
+ VerCapExtAttributeStorage(Attribute version, Attribute capabilities,
+ Attribute extensions)
+ : version(version), capabilities(capabilities), extensions(extensions) {}
+
+ bool operator==(const KeyTy &key) const {
+ return std::get<0>(key) == version && std::get<1>(key) == capabilities &&
+ std::get<2>(key) == extensions;
+ }
+
+ static VerCapExtAttributeStorage *
+ construct(AttributeStorageAllocator &allocator, const KeyTy &key) {
+ return new (allocator.allocate<VerCapExtAttributeStorage>())
+ VerCapExtAttributeStorage(std::get<0>(key), std::get<1>(key),
+ std::get<2>(key));
+ }
+
+ Attribute version;
+ Attribute capabilities;
+ Attribute extensions;
+};
+
struct TargetEnvAttributeStorage : public AttributeStorage {
- using KeyTy = std::tuple<Attribute, Attribute, Attribute, Attribute>;
+ using KeyTy = std::pair<Attribute, Attribute>;
- TargetEnvAttributeStorage(Attribute version, Attribute extensions,
- Attribute capabilities, Attribute limits)
- : version(version), extensions(extensions), capabilities(capabilities),
- limits(limits) {}
+ TargetEnvAttributeStorage(Attribute triple, Attribute limits)
+ : triple(triple), limits(limits) {}
bool operator==(const KeyTy &key) const {
- return std::get<0>(key) == version && std::get<1>(key) == extensions &&
- std::get<2>(key) == capabilities && std::get<3>(key) == limits;
+ return key.first == triple && key.second == limits;
}
static TargetEnvAttributeStorage *
construct(AttributeStorageAllocator &allocator, const KeyTy &key) {
return new (allocator.allocate<TargetEnvAttributeStorage>())
- TargetEnvAttributeStorage(std::get<0>(key), std::get<1>(key),
- std::get<2>(key), std::get<3>(key));
+ TargetEnvAttributeStorage(key.first, key.second);
}
- Attribute version;
- Attribute extensions;
- Attribute capabilities;
+ Attribute triple;
Attribute limits;
};
} // namespace detail
} // namespace spirv
} // namespace mlir
-spirv::TargetEnvAttr spirv::TargetEnvAttr::get(
- spirv::Version version, ArrayRef<spirv::Extension> extensions,
- ArrayRef<spirv::Capability> capabilities, DictionaryAttr limits) {
- Builder b(limits.getContext());
+//===----------------------------------------------------------------------===//
+// VerCapExtAttr
+//===----------------------------------------------------------------------===//
- auto versionAttr = b.getI32IntegerAttr(static_cast<uint32_t>(version));
+spirv::VerCapExtAttr spirv::VerCapExtAttr::get(
+ spirv::Version version, ArrayRef<spirv::Capability> capabilities,
+ ArrayRef<spirv::Extension> extensions, MLIRContext *context) {
+ Builder b(context);
- SmallVector<Attribute, 4> extAttrs;
- extAttrs.reserve(extensions.size());
- for (spirv::Extension ext : extensions)
- extAttrs.push_back(b.getStringAttr(spirv::stringifyExtension(ext)));
+ auto versionAttr = b.getI32IntegerAttr(static_cast<uint32_t>(version));
SmallVector<Attribute, 4> capAttrs;
capAttrs.reserve(capabilities.size());
for (spirv::Capability cap : capabilities)
capAttrs.push_back(b.getI32IntegerAttr(static_cast<uint32_t>(cap)));
- return get(versionAttr, b.getArrayAttr(extAttrs), b.getArrayAttr(capAttrs),
- limits);
+ SmallVector<Attribute, 4> extAttrs;
+ extAttrs.reserve(extensions.size());
+ for (spirv::Extension ext : extensions)
+ extAttrs.push_back(b.getStringAttr(spirv::stringifyExtension(ext)));
+
+ return get(versionAttr, b.getArrayAttr(capAttrs), b.getArrayAttr(extAttrs));
}
-spirv::TargetEnvAttr spirv::TargetEnvAttr::get(IntegerAttr version,
- ArrayAttr extensions,
+spirv::VerCapExtAttr spirv::VerCapExtAttr::get(IntegerAttr version,
ArrayAttr capabilities,
- DictionaryAttr limits) {
- assert(version && extensions && capabilities && limits);
+ ArrayAttr extensions) {
+ assert(version && capabilities && extensions);
MLIRContext *context = version.getContext();
- return Base::get(context, spirv::AttrKind::TargetEnv, version, extensions,
- capabilities, limits);
+ return Base::get(context, spirv::AttrKind::VerCapExt, version, capabilities,
+ extensions);
}
-StringRef spirv::TargetEnvAttr::getKindName() { return "target_env"; }
+StringRef spirv::VerCapExtAttr::getKindName() { return "vce"; }
-spirv::Version spirv::TargetEnvAttr::getVersion() {
+spirv::Version spirv::VerCapExtAttr::getVersion() {
return static_cast<spirv::Version>(
getImpl()->version.cast<IntegerAttr>().getValue().getZExtValue());
}
-spirv::TargetEnvAttr::ext_iterator::ext_iterator(ArrayAttr::iterator it)
+spirv::VerCapExtAttr::ext_iterator::ext_iterator(ArrayAttr::iterator it)
: llvm::mapped_iterator<ArrayAttr::iterator,
spirv::Extension (*)(Attribute)>(
it, [](Attribute attr) {
return *symbolizeExtension(attr.cast<StringAttr>().getValue());
}) {}
-spirv::TargetEnvAttr::ext_range spirv::TargetEnvAttr::getExtensions() {
+spirv::VerCapExtAttr::ext_range spirv::VerCapExtAttr::getExtensions() {
auto range = getExtensionsAttr().getValue();
return {ext_iterator(range.begin()), ext_iterator(range.end())};
}
-ArrayAttr spirv::TargetEnvAttr::getExtensionsAttr() {
+ArrayAttr spirv::VerCapExtAttr::getExtensionsAttr() {
return getImpl()->extensions.cast<ArrayAttr>();
}
-spirv::TargetEnvAttr::cap_iterator::cap_iterator(ArrayAttr::iterator it)
+spirv::VerCapExtAttr::cap_iterator::cap_iterator(ArrayAttr::iterator it)
: llvm::mapped_iterator<ArrayAttr::iterator,
spirv::Capability (*)(Attribute)>(
it, [](Attribute attr) {
@@ -110,25 +138,29 @@ spirv::TargetEnvAttr::cap_iterator::cap_iterator(ArrayAttr::iterator it)
attr.cast<IntegerAttr>().getValue().getZExtValue());
}) {}
-spirv::TargetEnvAttr::cap_range spirv::TargetEnvAttr::getCapabilities() {
+spirv::VerCapExtAttr::cap_range spirv::VerCapExtAttr::getCapabilities() {
auto range = getCapabilitiesAttr().getValue();
return {cap_iterator(range.begin()), cap_iterator(range.end())};
}
-ArrayAttr spirv::TargetEnvAttr::getCapabilitiesAttr() {
+ArrayAttr spirv::VerCapExtAttr::getCapabilitiesAttr() {
return getImpl()->capabilities.cast<ArrayAttr>();
}
-spirv::ResourceLimitsAttr spirv::TargetEnvAttr::getResourceLimits() {
- return getImpl()->limits.cast<spirv::ResourceLimitsAttr>();
-}
-
-LogicalResult spirv::TargetEnvAttr::verifyConstructionInvariants(
- Location loc, IntegerAttr version, ArrayAttr extensions,
- ArrayAttr capabilities, DictionaryAttr limits) {
+LogicalResult spirv::VerCapExtAttr::verifyConstructionInvariants(
+ Location loc, IntegerAttr version, ArrayAttr capabilities,
+ ArrayAttr extensions) {
if (!version.getType().isSignlessInteger(32))
return emitError(loc, "expected 32-bit integer for version");
+ if (!llvm::all_of(capabilities.getValue(), [](Attribute attr) {
+ if (auto intAttr = attr.dyn_cast<IntegerAttr>())
+ if (spirv::symbolizeCapability(intAttr.getValue().getZExtValue()))
+ return true;
+ return false;
+ }))
+ return emitError(loc, "unknown capability in capability list");
+
if (!llvm::all_of(extensions.getValue(), [](Attribute attr) {
if (auto strAttr = attr.dyn_cast<StringAttr>())
if (spirv::symbolizeExtension(strAttr.getValue()))
@@ -137,20 +169,62 @@ LogicalResult spirv::TargetEnvAttr::verifyConstructionInvariants(
}))
return emitError(loc, "unknown extension in extension list");
- if (!llvm::all_of(capabilities.getValue(), [](Attribute attr) {
- if (auto intAttr = attr.dyn_cast<IntegerAttr>())
- if (spirv::symbolizeCapability(intAttr.getValue().getZExtValue()))
- return true;
- return false;
- }))
- return emitError(loc, "unknown capability in capability list");
+ return success();
+}
+
+//===----------------------------------------------------------------------===//
+// TargetEnvAttr
+//===----------------------------------------------------------------------===//
+
+spirv::TargetEnvAttr spirv::TargetEnvAttr::get(spirv::VerCapExtAttr triple,
+ DictionaryAttr limits) {
+ assert(triple && limits && "expected valid triple and limits");
+ MLIRContext *context = triple.getContext();
+ return Base::get(context, spirv::AttrKind::TargetEnv, triple, limits);
+}
+
+StringRef spirv::TargetEnvAttr::getKindName() { return "target_env"; }
+spirv::VerCapExtAttr spirv::TargetEnvAttr::getTripleAttr() {
+ return getImpl()->triple.cast<spirv::VerCapExtAttr>();
+}
+
+spirv::Version spirv::TargetEnvAttr::getVersion() {
+ return getTripleAttr().getVersion();
+}
+
+spirv::VerCapExtAttr::ext_range spirv::TargetEnvAttr::getExtensions() {
+ return getTripleAttr().getExtensions();
+}
+
+ArrayAttr spirv::TargetEnvAttr::getExtensionsAttr() {
+ return getTripleAttr().getExtensionsAttr();
+}
+
+spirv::VerCapExtAttr::cap_range spirv::TargetEnvAttr::getCapabilities() {
+ return getTripleAttr().getCapabilities();
+}
+
+ArrayAttr spirv::TargetEnvAttr::getCapabilitiesAttr() {
+ return getTripleAttr().getCapabilitiesAttr();
+}
+
+spirv::ResourceLimitsAttr spirv::TargetEnvAttr::getResourceLimits() {
+ return getImpl()->limits.cast<spirv::ResourceLimitsAttr>();
+}
+
+LogicalResult spirv::TargetEnvAttr::verifyConstructionInvariants(
+ Location loc, spirv::VerCapExtAttr triple, DictionaryAttr limits) {
if (!limits.isa<spirv::ResourceLimitsAttr>())
return emitError(loc, "expected spirv::ResourceLimitsAttr for limits");
return success();
}
+//===----------------------------------------------------------------------===//
+// Utility functions
+//===----------------------------------------------------------------------===//
+
StringRef spirv::getInterfaceVarABIAttrName() {
return "spv.interface_var_abi";
}
@@ -212,13 +286,11 @@ spirv::getDefaultResourceLimits(MLIRContext *context) {
StringRef spirv::getTargetEnvAttrName() { return "spv.target_env"; }
spirv::TargetEnvAttr spirv::getDefaultTargetEnv(MLIRContext *context) {
- Builder builder(context);
- return spirv::TargetEnvAttr::get(
- builder.getI32IntegerAttr(static_cast<uint32_t>(spirv::Version::V_1_0)),
- builder.getI32ArrayAttr({}),
- builder.getI32ArrayAttr(
- {static_cast<uint32_t>(spirv::Capability::Shader)}),
- spirv::getDefaultResourceLimits(context));
+ auto triple = spirv::VerCapExtAttr::get(spirv::Version::V_1_0,
+ {spirv::Capability::Shader},
+ ArrayRef<Extension>(), context);
+ return spirv::TargetEnvAttr::get(triple,
+ spirv::getDefaultResourceLimits(context));
}
spirv::TargetEnvAttr spirv::lookupTargetEnvOrDefault(Operation *op) {
diff --git a/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir b/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir
index 6caaf8a00a36..cebd541977ef 100644
--- a/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir
+++ b/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir
@@ -16,7 +16,7 @@
module attributes {
spv.target_env = #spv.target_env<
- V_1_3, [], [Shader, GroupNonUniformArithmetic],
+ #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>,
{
max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>
@@ -78,7 +78,7 @@ func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>)
module attributes {
spv.target_env = #spv.target_env<
- V_1_3, [], [Shader, GroupNonUniformArithmetic],
+ #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>,
{
max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>
@@ -111,7 +111,7 @@ func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>)
module attributes {
spv.target_env = #spv.target_env<
- V_1_3, [], [Shader, GroupNonUniformArithmetic],
+ #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>,
{
max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>
@@ -146,7 +146,7 @@ func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>)
module attributes {
spv.target_env = #spv.target_env<
- V_1_3, [], [Shader, GroupNonUniformArithmetic],
+ #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>,
{
max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>
diff --git a/mlir/test/Dialect/SPIRV/availability.mlir b/mlir/test/Dialect/SPIRV/availability.mlir
index 381754c74609..a5203a0e4a2a 100644
--- a/mlir/test/Dialect/SPIRV/availability.mlir
+++ b/mlir/test/Dialect/SPIRV/availability.mlir
@@ -2,8 +2,8 @@
// CHECK-LABEL: iadd
func @iadd(%arg: i32) -> i32 {
- // CHECK: min version: V_1_0
- // CHECK: max version: V_1_5
+ // CHECK: min version: v1.0
+ // CHECK: max version: v1.5
// CHECK: extensions: [ ]
// CHECK: capabilities: [ ]
%0 = spv.IAdd %arg, %arg: i32
@@ -12,8 +12,8 @@ func @iadd(%arg: i32) -> i32 {
// CHECK: atomic_compare_exchange_weak
func @atomic_compare_exchange_weak(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 {
- // CHECK: min version: V_1_0
- // CHECK: max version: V_1_3
+ // CHECK: min version: v1.0
+ // CHECK: max version: v1.3
// CHECK: extensions: [ ]
// CHECK: capabilities: [ [Kernel] ]
%0 = spv.AtomicCompareExchangeWeak "Workgroup" "Release" "Acquire" %ptr, %value, %comparator: !spv.ptr<i32, Workgroup>
@@ -22,8 +22,8 @@ func @atomic_compare_exchange_weak(%ptr: !spv.ptr<i32, Workgroup>, %value: i32,
// CHECK-LABEL: subgroup_ballot
func @subgroup_ballot(%predicate: i1) -> vector<4xi32> {
- // CHECK: min version: V_1_3
- // CHECK: max version: V_1_5
+ // CHECK: min version: v1.3
+ // CHECK: max version: v1.5
// CHECK: extensions: [ ]
// CHECK: capabilities: [ [GroupNonUniformBallot] ]
%0 = spv.GroupNonUniformBallot "Workgroup" %predicate : vector<4xi32>
@@ -32,8 +32,8 @@ func @subgroup_ballot(%predicate: i1) -> vector<4xi32> {
// CHECK-LABEL: module_logical_glsl450
func @module_logical_glsl450() {
- // CHECK: spv.module min version: V_1_0
- // CHECK: spv.module max version: V_1_5
+ // CHECK: spv.module min version: v1.0
+ // CHECK: spv.module max version: v1.5
// CHECK: spv.module extensions: [ ]
// CHECK: spv.module capabilities: [ [Shader] ]
spv.module "Logical" "GLSL450" { }
@@ -42,8 +42,8 @@ func @module_logical_glsl450() {
// CHECK-LABEL: module_physical_storage_buffer64_vulkan
func @module_physical_storage_buffer64_vulkan() {
- // CHECK: spv.module min version: V_1_0
- // CHECK: spv.module max version: V_1_5
+ // CHECK: spv.module min version: v1.0
+ // CHECK: spv.module max version: v1.5
// CHECK: spv.module extensions: [ [SPV_EXT_physical_storage_buffer, SPV_KHR_physical_storage_buffer] [SPV_KHR_vulkan_memory_model] ]
// CHECK: spv.module capabilities: [ [PhysicalStorageBufferAddresses] [VulkanMemoryModel] ]
spv.module "PhysicalStorageBuffer64" "Vulkan" { }
diff --git a/mlir/test/Dialect/SPIRV/target-and-abi.mlir b/mlir/test/Dialect/SPIRV/target-and-abi.mlir
index 1182e28424e4..a28ca29e0ab9 100644
--- a/mlir/test/Dialect/SPIRV/target-and-abi.mlir
+++ b/mlir/test/Dialect/SPIRV/target-and-abi.mlir
@@ -106,87 +106,99 @@ func @interface_var() -> (f32 {spv.interface_var_abi = {
// spv.target_env
//===----------------------------------------------------------------------===//
-func @target_env_wrong_type() attributes {
- // expected-error @+1 {{expected valid keyword}}
- spv.target_env = #spv.target_env<64>
+func @target_env_missing_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_size = dense<[128, 64, 64]> : vector<3xi32>}>
} { return }
// -----
-func @target_env_missing_fields() attributes {
- // expected-error @+1 {{expected ','}}
- spv.target_env = #spv.target_env<V_1_0>
+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 @target_env_wrong_version() attributes {
- // expected-error @+1 {{unknown version: V_x_y}}
- spv.target_env = #spv.target_env<V_x_y, []>
+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>}>
+ 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>
+ }>
} { return }
// -----
-func @target_env_wrong_extension_type() attributes {
- // expected-error @+1 {{expected valid keyword}}
- spv.target_env = #spv.target_env<V_1_0, [32: i32], [Shader]>
+func @target_env_extra_fields() attributes {
+ // expected-error @+6 {{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>
+ },
+ more_stuff
+ >
} { return }
// -----
-func @target_env_wrong_extension() attributes {
- // expected-error @+1 {{unknown extension: SPV_Something}}
- spv.target_env = #spv.target_env<V_1_0, [SPV_Something], [Shader]>
+//===----------------------------------------------------------------------===//
+// spv.vce
+//===----------------------------------------------------------------------===//
+
+func @vce_wrong_type() attributes {
+ // expected-error @+1 {{expected valid keyword}}
+ vce = #spv.vce<64>
} { return }
// -----
-func @target_env_wrong_capability() attributes {
- // expected-error @+1 {{unknown capability: Something}}
- spv.target_env = #spv.target_env<V_1_0, [], [Something]>
+func @vce_missing_fields() attributes {
+ // expected-error @+1 {{expected ','}}
+ vce = #spv.vce<v1.0>
} { return }
// -----
-func @target_env_missing_limits() attributes {
- spv.target_env = #spv.target_env<
- V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader],
- // 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_size = dense<[128, 64, 64]> : vector<3xi32>}>
+func @vce_wrong_version() attributes {
+ // expected-error @+1 {{unknown version: V_x_y}}
+ vce = #spv.vce<V_x_y, []>
} { return }
// -----
-func @target_env_wrong_limits() attributes {
- spv.target_env = #spv.target_env<
- V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader],
- // 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>}>
+func @vce_wrong_extension_type() attributes {
+ // expected-error @+1 {{expected valid keyword}}
+ vce = #spv.vce<v1.0, [32: i32], [Shader]>
} { return }
// -----
-func @target_env() attributes {
+func @vce_wrong_extension() attributes {
+ // expected-error @+1 {{unknown extension: SPV_Something}}
+ vce = #spv.vce<v1.0, [Shader], [SPV_Something]>
+} { return }
- // CHECK: spv.target_env = #spv.target_env<V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
- spv.target_env = #spv.target_env<
- V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader],
- {
- max_compute_workgroup_invocations = 128 : i32,
- max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>
- }>
+// -----
+
+func @vce_wrong_capability() attributes {
+ // expected-error @+1 {{unknown capability: Something}}
+ vce = #spv.vce<v1.0, [Something], []>
} { return }
// -----
-func @target_env_extra_fields() attributes {
- // expected-error @+6 {{expected '>'}}
- spv.target_env = #spv.target_env<
- V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader],
- {
- max_compute_workgroup_invocations = 128 : i32,
- max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>
- },
- more_stuff
- >
+func @vce() attributes {
+ // CHECK: #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>
+ vce = #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>
} { return }
diff --git a/mlir/test/Dialect/SPIRV/target-env.mlir b/mlir/test/Dialect/SPIRV/target-env.mlir
index 1e43ec9fbb9d..32f36e96f5ea 100644
--- a/mlir/test/Dialect/SPIRV/target-env.mlir
+++ b/mlir/test/Dialect/SPIRV/target-env.mlir
@@ -35,7 +35,7 @@
// CHECK-LABEL: @cmp_exchange_weak_suitable_version_capabilities
func @cmp_exchange_weak_suitable_version_capabilities(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
- spv.target_env = #spv.target_env<V_1_1, [], [Kernel, AtomicStorage], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.1, [Kernel, AtomicStorage], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// 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 @cmp_exchange_weak_suitable_version_capabilities(%ptr: !spv.ptr<i32, Workgr
// CHECK-LABEL: @cmp_exchange_weak_unsupported_version
func @cmp_exchange_weak_unsupported_version(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
- spv.target_env = #spv.target_env<V_1_4, [], [Kernel, AtomicStorage], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.4, [Kernel, AtomicStorage], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// 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 @cmp_exchange_weak_unsupported_version(%ptr: !spv.ptr<i32, Workgroup>, %val
// CHECK-LABEL: @group_non_uniform_ballot_suitable_version
func @group_non_uniform_ballot_suitable_version(%predicate: i1) -> vector<4xi32> attributes {
- spv.target_env = #spv.target_env<V_1_4, [], [GroupNonUniformBallot], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.4, [GroupNonUniformBallot], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// CHECK: spv.GroupNonUniformBallot "Workgroup"
%0 = "test.convert_to_group_non_uniform_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
@@ -66,7 +66,7 @@ func @group_non_uniform_ballot_suitable_version(%predicate: i1) -> vector<4xi32>
// CHECK-LABEL: @group_non_uniform_ballot_unsupported_version
func @group_non_uniform_ballot_unsupported_version(%predicate: i1) -> vector<4xi32> attributes {
- spv.target_env = #spv.target_env<V_1_1, [], [GroupNonUniformBallot], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.1, [GroupNonUniformBallot], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// 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 @group_non_uniform_ballot_unsupported_version(%predicate: i1) -> vector<4xi
// CHECK-LABEL: @cmp_exchange_weak_missing_capability_kernel
func @cmp_exchange_weak_missing_capability_kernel(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
- spv.target_env = #spv.target_env<V_1_3, [], [AtomicStorage], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.3, [AtomicStorage], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// 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 @cmp_exchange_weak_missing_capability_kernel(%ptr: !spv.ptr<i32, Workgroup>
// CHECK-LABEL: @cmp_exchange_weak_missing_capability_atomic_storage
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<V_1_3, [], [Kernel], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.3, [Kernel], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// 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 @cmp_exchange_weak_missing_capability_atomic_storage(%ptr: !spv.ptr<i32, Wo
// CHECK-LABEL: @subgroup_ballot_missing_capability
func @subgroup_ballot_missing_capability(%predicate: i1) -> vector<4xi32> attributes {
- spv.target_env = #spv.target_env<V_1_4, [SPV_KHR_shader_ballot], [], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.4, [], [SPV_KHR_shader_ballot]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// CHECK: test.convert_to_subgroup_ballot_op
%0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
@@ -106,7 +106,7 @@ func @subgroup_ballot_missing_capability(%predicate: i1) -> vector<4xi32> attrib
// CHECK-LABEL: @bit_reverse_directly_implied_capability
func @bit_reverse_directly_implied_capability(%operand: i32) -> i32 attributes {
- spv.target_env = #spv.target_env<V_1_0, [], [Geometry], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.0, [Geometry], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// CHECK: spv.BitReverse
%0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32)
@@ -115,7 +115,7 @@ func @bit_reverse_directly_implied_capability(%operand: i32) -> i32 attributes {
// CHECK-LABEL: @bit_reverse_recursively_implied_capability
func @bit_reverse_recursively_implied_capability(%operand: i32) -> i32 attributes {
- spv.target_env = #spv.target_env<V_1_0, [], [GeometryPointSize], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.0, [GeometryPointSize], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// CHECK: spv.BitReverse
%0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32)
@@ -128,7 +128,7 @@ func @bit_reverse_recursively_implied_capability(%operand: i32) -> i32 attribute
// CHECK-LABEL: @subgroup_ballot_suitable_extension
func @subgroup_ballot_suitable_extension(%predicate: i1) -> vector<4xi32> attributes {
- spv.target_env = #spv.target_env<V_1_4, [SPV_KHR_shader_ballot], [SubgroupBallotKHR], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], [SPV_KHR_shader_ballot]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// CHECK: spv.SubgroupBallotKHR
%0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
@@ -137,7 +137,7 @@ func @subgroup_ballot_suitable_extension(%predicate: i1) -> vector<4xi32> attrib
// CHECK-LABEL: @subgroup_ballot_missing_extension
func @subgroup_ballot_missing_extension(%predicate: i1) -> vector<4xi32> attributes {
- spv.target_env = #spv.target_env<V_1_4, [], [SubgroupBallotKHR], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// CHECK: test.convert_to_subgroup_ballot_op
%0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
@@ -146,7 +146,7 @@ func @subgroup_ballot_missing_extension(%predicate: i1) -> vector<4xi32> attribu
// CHECK-LABEL: @module_suitable_extension1
func @module_suitable_extension1() attributes {
- spv.target_env = #spv.target_env<V_1_0, [SPV_KHR_vulkan_memory_model, SPV_EXT_physical_storage_buffer], [VulkanMemoryModel, PhysicalStorageBufferAddresses], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_EXT_physical_storage_buffer]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// CHECK: spv.module "PhysicalStorageBuffer64" "Vulkan"
"test.convert_to_module_op"() : () ->()
@@ -155,7 +155,7 @@ func @module_suitable_extension1() attributes {
// CHECK-LABEL: @module_suitable_extension2
func @module_suitable_extension2() attributes {
- spv.target_env = #spv.target_env<V_1_0, [SPV_KHR_vulkan_memory_model, SPV_KHR_physical_storage_buffer], [VulkanMemoryModel, PhysicalStorageBufferAddresses], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_KHR_physical_storage_buffer]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// CHECK: spv.module "PhysicalStorageBuffer64" "Vulkan"
"test.convert_to_module_op"() : () -> ()
@@ -164,7 +164,7 @@ func @module_suitable_extension2() attributes {
// CHECK-LABEL: @module_missing_extension_mm
func @module_missing_extension_mm() attributes {
- spv.target_env = #spv.target_env<V_1_0, [SPV_KHR_physical_storage_buffer], [VulkanMemoryModel, PhysicalStorageBufferAddresses], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_physical_storage_buffer]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// CHECK: test.convert_to_module_op
"test.convert_to_module_op"() : () -> ()
@@ -173,7 +173,7 @@ func @module_missing_extension_mm() attributes {
// CHECK-LABEL: @module_missing_extension_am
func @module_missing_extension_am() attributes {
- spv.target_env = #spv.target_env<V_1_0, [SPV_KHR_vulkan_memory_model], [VulkanMemoryModel, PhysicalStorageBufferAddresses], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// CHECK: test.convert_to_module_op
"test.convert_to_module_op"() : () -> ()
@@ -183,7 +183,7 @@ func @module_missing_extension_am() attributes {
// CHECK-LABEL: @module_implied_extension
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<V_1_5, [], [VulkanMemoryModel, PhysicalStorageBufferAddresses], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+ spv.target_env = #spv.target_env<#spv.vce<v1.5, [VulkanMemoryModel, PhysicalStorageBufferAddresses], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
// CHECK: spv.module "PhysicalStorageBuffer64" "Vulkan"
"test.convert_to_module_op"() : () -> ()
More information about the Mlir-commits
mailing list