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` | ...
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`
```
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>
// 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]>;
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
};
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
};
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();
static bool kindof(unsigned kind) { return kind == AttrKind::TargetEnv; }
static LogicalResult verifyConstructionInvariants(Location loc,
- IntegerAttr version,
- ArrayAttr extensions,
- ArrayAttr capabilities,
+ VerCapExtAttr triple,
DictionaryAttr limits);
};
: Dialect(getDialectNamespace(), context) {
addTypes<ArrayType, ImageType, PointerType, RuntimeArrayType, StructType>();
- addAttributes<TargetEnvAttr>();
+ addAttributes<TargetEnvAttr, VerCapExtAttr>();
// Add SPIR-V ops.
addOperations<
return success();
}
-/// Parses a spirv::TargetEnvAttr.
-static Attribute parseTargetAttr(DialectAsmParser &parser) {
+static Attribute parseVerCapExtAttr(DialectAsmParser &parser) {
if (parser.parseLess())
return {};
}
}
- 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();
if (parser.parseGreater())
return {};
- return spirv::TargetEnvAttr::get(versionAttr, extensionsAttr,
- capabilitiesAttr, limitsAttr);
+ return spirv::TargetEnvAttr::get(tripleAttr, limitsAttr);
}
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;
// 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");
}
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);
"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";
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) {
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()))
}))
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";
}
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) {
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>
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>
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>
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>
// 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
// 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>
// 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>
// 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" { }
// 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" { }
// 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 }
// 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)
// 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)
// 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>)
// 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>)
// 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)
// 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)
// 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>)
// 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)
// 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)
// 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>)
// 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>)
// 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"() : () ->()
// 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"() : () -> ()
// 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"() : () -> ()
// 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"() : () -> ()
// 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"() : () -> ()
projects / platform / upstream / llvm.git / commitdiff