forked from OSchip/llvm-project
[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
This commit is contained in:
Lei Zhang
parent
69ce2fd2df
commit
e115a40f50
|
|
@ -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>
|
||||
|
|
|
|||
|
|
@ -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]>;
|
||||
|
|
|
|||
|
|
@ -32,16 +32,71 @@ 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 SPIR-V (version, capabilities, extensions)
|
||||
/// triple.
|
||||
class VerCapExtAttr
|
||||
: public Attribute::AttrBase<VerCapExtAttr, Attribute,
|
||||
detail::VerCapExtAttributeStorage> {
|
||||
public:
|
||||
using Base::Base;
|
||||
|
||||
/// 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 version.
|
||||
Version getVersion();
|
||||
|
||||
struct ext_iterator final
|
||||
: public llvm::mapped_iterator<ArrayAttr::iterator,
|
||||
Extension (*)(Attribute)> {
|
||||
explicit ext_iterator(ArrayAttr::iterator it);
|
||||
};
|
||||
using ext_range = llvm::iterator_range<ext_iterator>;
|
||||
|
||||
/// Returns the extensions.
|
||||
ext_range getExtensions();
|
||||
/// Returns the extensions as a string array attribute.
|
||||
ArrayAttr getExtensionsAttr();
|
||||
|
||||
struct cap_iterator final
|
||||
: public llvm::mapped_iterator<ArrayAttr::iterator,
|
||||
Capability (*)(Attribute)> {
|
||||
explicit cap_iterator(ArrayAttr::iterator it);
|
||||
};
|
||||
using cap_range = llvm::iterator_range<cap_iterator>;
|
||||
|
||||
/// 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.
|
||||
|
|
@ -52,39 +107,24 @@ 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);
|
||||
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();
|
||||
|
||||
struct ext_iterator final
|
||||
: public llvm::mapped_iterator<ArrayAttr::iterator,
|
||||
Extension (*)(Attribute)> {
|
||||
explicit ext_iterator(ArrayAttr::iterator it);
|
||||
};
|
||||
using ext_range = llvm::iterator_range<ext_iterator>;
|
||||
|
||||
/// Returns the target extensions.
|
||||
ext_range getExtensions();
|
||||
VerCapExtAttr::ext_range getExtensions();
|
||||
/// Returns the target extensions as a string array attribute.
|
||||
ArrayAttr getExtensionsAttr();
|
||||
|
||||
struct cap_iterator final
|
||||
: public llvm::mapped_iterator<ArrayAttr::iterator,
|
||||
Capability (*)(Attribute)> {
|
||||
explicit cap_iterator(ArrayAttr::iterator it);
|
||||
};
|
||||
using cap_range = llvm::iterator_range<cap_iterator>;
|
||||
|
||||
/// Returns the target capabilities.
|
||||
cap_range getCapabilities();
|
||||
VerCapExtAttr::cap_range getCapabilities();
|
||||
/// Returns the target capabilities as an integer array attribute.
|
||||
ArrayAttr getCapabilitiesAttr();
|
||||
|
||||
|
|
@ -94,9 +134,7 @@ public:
|
|||
static bool kindof(unsigned kind) { return kind == AttrKind::TargetEnv; }
|
||||
|
||||
static LogicalResult verifyConstructionInvariants(Location loc,
|
||||
IntegerAttr version,
|
||||
ArrayAttr extensions,
|
||||
ArrayAttr capabilities,
|
||||
VerCapExtAttr triple,
|
||||
DictionaryAttr limits);
|
||||
};
|
||||
|
||||
|
|
|
|||
|
|
@ -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,28 +684,6 @@ static Attribute parseTargetAttr(DialectAsmParser &parser) {
|
|||
}
|
||||
}
|
||||
|
||||
ArrayAttr extensionsAttr;
|
||||
{
|
||||
SmallVector<Attribute, 1> extensions;
|
||||
llvm::SMLoc errorloc;
|
||||
StringRef errorKeyword;
|
||||
|
||||
auto processExtension = [&](llvm::SMLoc loc, StringRef extension) {
|
||||
if (spirv::symbolizeExtension(extension)) {
|
||||
extensions.push_back(builder.getStringAttr(extension));
|
||||
return success();
|
||||
}
|
||||
return errorloc = loc, errorKeyword = extension, failure();
|
||||
};
|
||||
if (parseKeywordList(parser, processExtension) || parser.parseComma()) {
|
||||
if (!errorKeyword.empty())
|
||||
parser.emitError(errorloc, "unknown extension: ") << errorKeyword;
|
||||
return {};
|
||||
}
|
||||
|
||||
extensionsAttr = builder.getArrayAttr(extensions);
|
||||
}
|
||||
|
||||
ArrayAttr capabilitiesAttr;
|
||||
{
|
||||
SmallVector<Attribute, 4> capabilities;
|
||||
|
|
@ -730,6 +707,44 @@ static Attribute parseTargetAttr(DialectAsmParser &parser) {
|
|||
capabilitiesAttr = builder.getArrayAttr(capabilities);
|
||||
}
|
||||
|
||||
ArrayAttr extensionsAttr;
|
||||
{
|
||||
SmallVector<Attribute, 1> extensions;
|
||||
llvm::SMLoc errorloc;
|
||||
StringRef errorKeyword;
|
||||
|
||||
auto processExtension = [&](llvm::SMLoc loc, StringRef extension) {
|
||||
if (spirv::symbolizeExtension(extension)) {
|
||||
extensions.push_back(builder.getStringAttr(extension));
|
||||
return success();
|
||||
}
|
||||
return errorloc = loc, errorKeyword = extension, failure();
|
||||
};
|
||||
if (parseKeywordList(parser, processExtension)) {
|
||||
if (!errorKeyword.empty())
|
||||
parser.emitError(errorloc, "unknown extension: ") << errorKeyword;
|
||||
return {};
|
||||
}
|
||||
|
||||
extensionsAttr = builder.getArrayAttr(extensions);
|
||||
}
|
||||
|
||||
if (parser.parseGreater())
|
||||
return {};
|
||||
|
||||
return spirv::VerCapExtAttr::get(versionAttr, capabilitiesAttr,
|
||||
extensionsAttr);
|
||||
}
|
||||
|
||||
/// Parses a spirv::TargetEnvAttr.
|
||||
static Attribute parseTargetEnvAttr(DialectAsmParser &parser) {
|
||||
if (parser.parseLess())
|
||||
return {};
|
||||
|
||||
spirv::VerCapExtAttr tripleAttr;
|
||||
if (parser.parseAttribute(tripleAttr) || parser.parseComma())
|
||||
return {};
|
||||
|
||||
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 << "], [";
|
||||
interleaveComma(targetEnv.getCapabilities(), os, [&](spirv::Capability cap) {
|
||||
printer << spirv::VerCapExtAttr::getKindName() << "<"
|
||||
<< spirv::stringifyVersion(triple.getVersion()) << ", [";
|
||||
interleaveComma(triple.getCapabilities(), os, [&](spirv::Capability cap) {
|
||||
os << spirv::stringifyCapability(cap);
|
||||
});
|
||||
printer << "], " << targetEnv.getResourceLimits() << ">";
|
||||
printer << "], [";
|
||||
interleaveComma(triple.getExtensionsAttr(), os, [&](Attribute attr) {
|
||||
os << attr.cast<StringAttr>().getValue();
|
||||
});
|
||||
printer << "]>";
|
||||
}
|
||||
|
||||
static void print(spirv::TargetEnvAttr targetEnv, DialectAsmPrinter &printer) {
|
||||
printer << spirv::TargetEnvAttr::getKindName() << "<#spv.";
|
||||
print(targetEnv.getTripleAttr(), printer);
|
||||
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";
|
||||
|
|
|
|||
|
|
@ -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 TargetEnvAttributeStorage : public AttributeStorage {
|
||||
using KeyTy = std::tuple<Attribute, Attribute, Attribute, Attribute>;
|
||||
struct VerCapExtAttributeStorage : public AttributeStorage {
|
||||
using KeyTy = std::tuple<Attribute, Attribute, Attribute>;
|
||||
|
||||
TargetEnvAttributeStorage(Attribute version, Attribute extensions,
|
||||
Attribute capabilities, Attribute limits)
|
||||
: version(version), extensions(extensions), capabilities(capabilities),
|
||||
limits(limits) {}
|
||||
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) == extensions &&
|
||||
std::get<2>(key) == capabilities && std::get<3>(key) == limits;
|
||||
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::pair<Attribute, Attribute>;
|
||||
|
||||
TargetEnvAttributeStorage(Attribute triple, Attribute limits)
|
||||
: triple(triple), limits(limits) {}
|
||||
|
||||
bool operator==(const KeyTy &key) const {
|
||||
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
|
||||
//===----------------------------------------------------------------------===//
|
||||
|
||||
spirv::VerCapExtAttr spirv::VerCapExtAttr::get(
|
||||
spirv::Version version, ArrayRef<spirv::Capability> capabilities,
|
||||
ArrayRef<spirv::Extension> extensions, MLIRContext *context) {
|
||||
Builder b(context);
|
||||
|
||||
auto versionAttr = b.getI32IntegerAttr(static_cast<uint32_t>(version));
|
||||
|
||||
SmallVector<Attribute, 4> extAttrs;
|
||||
extAttrs.reserve(extensions.size());
|
||||
for (spirv::Extension ext : extensions)
|
||||
extAttrs.push_back(b.getStringAttr(spirv::stringifyExtension(ext)));
|
||||
|
||||
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,33 +138,21 @@ 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(extensions.getValue(), [](Attribute attr) {
|
||||
if (auto strAttr = attr.dyn_cast<StringAttr>())
|
||||
if (spirv::symbolizeExtension(strAttr.getValue()))
|
||||
return true;
|
||||
return false;
|
||||
}))
|
||||
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()))
|
||||
|
|
@ -145,12 +161,70 @@ LogicalResult spirv::TargetEnvAttr::verifyConstructionInvariants(
|
|||
}))
|
||||
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 true;
|
||||
return false;
|
||||
}))
|
||||
return emitError(loc, "unknown extension in extension 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) {
|
||||
|
|
|
|||
|
|
@ -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>
|
||||
|
|
|
|||
|
|
@ -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" { }
|
||||
|
|
|
|||
|
|
@ -106,51 +106,9 @@ 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>
|
||||
} { return }
|
||||
|
||||
// -----
|
||||
|
||||
func @target_env_missing_fields() attributes {
|
||||
// expected-error @+1 {{expected ','}}
|
||||
spv.target_env = #spv.target_env<V_1_0>
|
||||
} { return }
|
||||
|
||||
// -----
|
||||
|
||||
func @target_env_wrong_version() attributes {
|
||||
// expected-error @+1 {{unknown version: V_x_y}}
|
||||
spv.target_env = #spv.target_env<V_x_y, []>
|
||||
} { 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]>
|
||||
} { 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]>
|
||||
} { return }
|
||||
|
||||
// -----
|
||||
|
||||
func @target_env_wrong_capability() attributes {
|
||||
// expected-error @+1 {{unknown capability: Something}}
|
||||
spv.target_env = #spv.target_env<V_1_0, [], [Something]>
|
||||
} { return }
|
||||
|
||||
// -----
|
||||
|
||||
func @target_env_missing_limits() attributes {
|
||||
spv.target_env = #spv.target_env<
|
||||
V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader],
|
||||
#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 }
|
||||
|
|
@ -159,7 +117,7 @@ func @target_env_missing_limits() attributes {
|
|||
|
||||
func @target_env_wrong_limits() attributes {
|
||||
spv.target_env = #spv.target_env<
|
||||
V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader],
|
||||
#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 }
|
||||
|
|
@ -167,10 +125,11 @@ func @target_env_wrong_limits() attributes {
|
|||
// -----
|
||||
|
||||
func @target_env() attributes {
|
||||
|
||||
// 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>}>
|
||||
// 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<
|
||||
V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader],
|
||||
#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>
|
||||
|
|
@ -182,7 +141,7 @@ func @target_env() attributes {
|
|||
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],
|
||||
#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>
|
||||
|
|
@ -190,3 +149,56 @@ func @target_env_extra_fields() attributes {
|
|||
more_stuff
|
||||
>
|
||||
} { return }
|
||||
|
||||
// -----
|
||||
|
||||
//===----------------------------------------------------------------------===//
|
||||
// spv.vce
|
||||
//===----------------------------------------------------------------------===//
|
||||
|
||||
func @vce_wrong_type() attributes {
|
||||
// expected-error @+1 {{expected valid keyword}}
|
||||
vce = #spv.vce<64>
|
||||
} { return }
|
||||
|
||||
// -----
|
||||
|
||||
func @vce_missing_fields() attributes {
|
||||
// expected-error @+1 {{expected ','}}
|
||||
vce = #spv.vce<v1.0>
|
||||
} { return }
|
||||
|
||||
// -----
|
||||
|
||||
func @vce_wrong_version() attributes {
|
||||
// expected-error @+1 {{unknown version: V_x_y}}
|
||||
vce = #spv.vce<V_x_y, []>
|
||||
} { return }
|
||||
|
||||
// -----
|
||||
|
||||
func @vce_wrong_extension_type() attributes {
|
||||
// expected-error @+1 {{expected valid keyword}}
|
||||
vce = #spv.vce<v1.0, [32: i32], [Shader]>
|
||||
} { return }
|
||||
|
||||
// -----
|
||||
|
||||
func @vce_wrong_extension() attributes {
|
||||
// expected-error @+1 {{unknown extension: SPV_Something}}
|
||||
vce = #spv.vce<v1.0, [Shader], [SPV_Something]>
|
||||
} { return }
|
||||
|
||||
// -----
|
||||
|
||||
func @vce_wrong_capability() attributes {
|
||||
// expected-error @+1 {{unknown capability: Something}}
|
||||
vce = #spv.vce<v1.0, [Something], []>
|
||||
} { return }
|
||||
|
||||
// -----
|
||||
|
||||
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 }
|
||||
|
|
|
|||
|
|
@ -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"() : () -> ()
|
||||
|
|
|
|||
Loading…
Reference in New Issue