diff --git a/mlir/docs/Dialects/SPIR-V.md b/mlir/docs/Dialects/SPIR-V.md index d81d1038ca4351..4380c9f006d1ce 100644 --- a/mlir/docs/Dialects/SPIR-V.md +++ b/mlir/docs/Dialects/SPIR-V.md @@ -742,11 +742,11 @@ instructions. SPIR-V compilation should also take into consideration of the execution environment, so we generate SPIR-V modules valid for the target environment. -This is conveyed by the `spv.target_env` attribute. It should be of -`#spv.target_env` attribute kind, which is defined as: +This is conveyed by the `spv.target_env` (`spirv::TargetEnvAttr`) attribute. It +should be of `#spv.target_env` attribute kind, which is defined as: ``` -spirv-version ::= `V_1_0` | `V_1_1` | ... +spirv-version ::= `v1.0` | `v1.1` | ... spirv-extension ::= `SPV_KHR_16bit_storage` | `SPV_EXT_physical_storage_buffer` | ... spirv-capability ::= `Shader` | `Kernel` | `GroupNonUniform` | ... @@ -758,18 +758,22 @@ spirv-capability-elements ::= spirv-capability (`,` spirv-capability)* spirv-resource-limits ::= dictionary-attribute +spirv-vce-attribute ::= `#` `spv.vce` `<` + spirv-version `,` + spirv-capability-list `,` + spirv-extensions-list `>` + spirv-target-env-attribute ::= `#` `spv.target_env` `<` - spirv-version `,` - spirv-extensions-list `,` - spirv-capability-list `,` + spirv-vce-attribute, spirv-resource-limits `>` ``` The attribute has a few fields: -* The target SPIR-V version. -* A list of SPIR-V extensions for the target. -* A list of SPIR-V capabilities for the target. +* A `#spv.vce` (`spirv::VerCapExtAttr`) attribute: + * The target SPIR-V version. + * A list of SPIR-V extensions for the target. + * A list of SPIR-V capabilities for the target. * A dictionary of target resource limits (see the [Vulkan spec][VulkanResourceLimits] for explanation): * `max_compute_workgroup_invocations` @@ -780,7 +784,7 @@ For example, ``` module attributes { spv.target_env = #spv.target_env< - V_1_3, [SPV_KHR_8bit_storage], [Shader, GroupNonUniform] + #spv.vce, { max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32> diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td index 4eefc618990345..1af6ddef4ea042 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td @@ -96,12 +96,12 @@ class SPV_StrEnumAttr; -def SPV_V_1_1 : I32EnumAttrCase<"V_1_1", 1>; -def SPV_V_1_2 : I32EnumAttrCase<"V_1_2", 2>; -def SPV_V_1_3 : I32EnumAttrCase<"V_1_3", 3>; -def SPV_V_1_4 : I32EnumAttrCase<"V_1_4", 4>; -def SPV_V_1_5 : I32EnumAttrCase<"V_1_5", 5>; +def SPV_V_1_0 : I32EnumAttrCase<"V_1_0", 0, "v1.0">; +def SPV_V_1_1 : I32EnumAttrCase<"V_1_1", 1, "v1.1">; +def SPV_V_1_2 : I32EnumAttrCase<"V_1_2", 2, "v1.2">; +def SPV_V_1_3 : I32EnumAttrCase<"V_1_3", 3, "v1.3">; +def SPV_V_1_4 : I32EnumAttrCase<"V_1_4", 4, "v1.4">; +def SPV_V_1_5 : I32EnumAttrCase<"V_1_5", 5, "v1.5">; def SPV_VersionAttr : SPV_I32EnumAttr<"Version", "valid SPIR-V version", [ SPV_V_1_0, SPV_V_1_1, SPV_V_1_2, SPV_V_1_3, SPV_V_1_4, SPV_V_1_5]>; diff --git a/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h b/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h index 01b775846ee312..1d3964a67fe54b 100644 --- a/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h +++ b/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h @@ -32,36 +32,37 @@ enum class Version : uint32_t; namespace detail { struct TargetEnvAttributeStorage; +struct VerCapExtAttributeStorage; } // namespace detail /// SPIR-V dialect-specific attribute kinds. // TODO(antiagainst): move to a more suitable place if we have more attributes. namespace AttrKind { enum Kind { - TargetEnv = Attribute::FIRST_SPIRV_ATTR, + TargetEnv = Attribute::FIRST_SPIRV_ATTR, /// Target environment + VerCapExt, /// (version, extension, capability) triple }; } // namespace AttrKind -/// An attribute that specifies the target version, allowed extensions and -/// capabilities, and resource limits. These information describles a SPIR-V -/// target environment. -class TargetEnvAttr - : public Attribute::AttrBase { +/// An attribute that specifies the SPIR-V (version, capabilities, extensions) +/// triple. +class VerCapExtAttr + : public Attribute::AttrBase { public: using Base::Base; - /// Gets a TargetEnvAttr instance. - static TargetEnvAttr get(Version version, ArrayRef extensions, - ArrayRef capabilities, - DictionaryAttr limits); - static TargetEnvAttr get(IntegerAttr version, ArrayAttr extensions, - ArrayAttr capabilities, DictionaryAttr limits); + /// Gets a VerCapExtAttr instance. + static VerCapExtAttr get(Version version, ArrayRef capabilities, + ArrayRef extensions, + MLIRContext *context); + static VerCapExtAttr get(IntegerAttr version, ArrayAttr capabilities, + ArrayAttr extensions); /// Returns the attribute kind's name (without the 'spv.' prefix). static StringRef getKindName(); - /// Returns the target version. + /// Returns the version. Version getVersion(); struct ext_iterator final @@ -71,9 +72,9 @@ class TargetEnvAttr }; using ext_range = llvm::iterator_range; - /// Returns the target extensions. + /// Returns the extensions. ext_range getExtensions(); - /// Returns the target extensions as a string array attribute. + /// Returns the extensions as a string array attribute. ArrayAttr getExtensionsAttr(); struct cap_iterator final @@ -83,8 +84,47 @@ class TargetEnvAttr }; using cap_range = llvm::iterator_range; - /// 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 { +public: + using Base::Base; + + /// Gets a TargetEnvAttr instance. + static TargetEnvAttr get(VerCapExtAttr triple, DictionaryAttr limits); + + /// Returns the attribute kind's name (without the 'spv.' prefix). + static StringRef getKindName(); + + /// Returns the (version, capabilities, extensions) triple attribute. + VerCapExtAttr getTripleAttr(); + + /// Returns the target version. + Version getVersion(); + + /// Returns the target extensions. + VerCapExtAttr::ext_range getExtensions(); + /// Returns the target extensions as a string array attribute. + ArrayAttr getExtensionsAttr(); + + /// Returns the target capabilities. + VerCapExtAttr::cap_range getCapabilities(); /// Returns the target capabilities as an integer array attribute. ArrayAttr getCapabilitiesAttr(); @@ -94,9 +134,7 @@ class TargetEnvAttr static bool kindof(unsigned kind) { return kind == AttrKind::TargetEnv; } static LogicalResult verifyConstructionInvariants(Location loc, - IntegerAttr version, - ArrayAttr extensions, - ArrayAttr capabilities, + VerCapExtAttr triple, DictionaryAttr limits); }; diff --git a/mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp b/mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp index 1946bfc37ce350..50ecf9ef7cbdaf 100644 --- a/mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp +++ b/mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp @@ -118,7 +118,7 @@ SPIRVDialect::SPIRVDialect(MLIRContext *context) : Dialect(getDialectNamespace(), context) { addTypes(); - addAttributes(); + addAttributes(); // Add SPIR-V ops. addOperations< @@ -662,8 +662,7 @@ static ParseResult parseKeywordList( return success(); } -/// Parses a spirv::TargetEnvAttr. -static Attribute parseTargetAttr(DialectAsmParser &parser) { +static Attribute parseVerCapExtAttr(DialectAsmParser &parser) { if (parser.parseLess()) return {}; @@ -685,51 +684,67 @@ static Attribute parseTargetAttr(DialectAsmParser &parser) { } } - ArrayAttr extensionsAttr; + ArrayAttr capabilitiesAttr; { - SmallVector extensions; + SmallVector 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(*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 capabilities; + SmallVector 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(*capSymbol))); + auto processExtension = [&](llvm::SMLoc loc, StringRef extension) { + if (spirv::symbolizeExtension(extension)) { + extensions.push_back(builder.getStringAttr(extension)); return success(); } - return errorloc = loc, errorKeyword = capability, failure(); + return errorloc = loc, errorKeyword = extension, failure(); }; - if (parseKeywordList(parser, processCapability) || parser.parseComma()) { + if (parseKeywordList(parser, processExtension)) { if (!errorKeyword.empty()) - parser.emitError(errorloc, "unknown capability: ") << errorKeyword; + parser.emitError(errorloc, "unknown extension: ") << errorKeyword; return {}; } - capabilitiesAttr = builder.getArrayAttr(capabilities); + extensionsAttr = builder.getArrayAttr(extensions); } + if (parser.parseGreater()) + return {}; + + return spirv::VerCapExtAttr::get(versionAttr, capabilitiesAttr, + extensionsAttr); +} + +/// Parses a spirv::TargetEnvAttr. +static Attribute parseTargetEnvAttr(DialectAsmParser &parser) { + if (parser.parseLess()) + return {}; + + spirv::VerCapExtAttr tripleAttr; + if (parser.parseAttribute(tripleAttr) || parser.parseComma()) + return {}; + DictionaryAttr limitsAttr; { auto loc = parser.getCurrentLocation(); @@ -749,8 +764,7 @@ static Attribute parseTargetAttr(DialectAsmParser &parser) { if (parser.parseGreater()) return {}; - return spirv::TargetEnvAttr::get(versionAttr, extensionsAttr, - capabilitiesAttr, limitsAttr); + return spirv::TargetEnvAttr::get(tripleAttr, limitsAttr); } Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser, @@ -767,7 +781,9 @@ Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser, return {}; if (attrKind == spirv::TargetEnvAttr::getKindName()) - return parseTargetAttr(parser); + return parseTargetEnvAttr(parser); + if (attrKind == spirv::VerCapExtAttr::getKindName()) + return parseVerCapExtAttr(parser); parser.emitError(parser.getNameLoc(), "unknown SPIR-V attriubte kind: ") << attrKind; @@ -778,24 +794,32 @@ Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser, // Attribute Printing //===----------------------------------------------------------------------===// -static void print(spirv::TargetEnvAttr targetEnv, DialectAsmPrinter &printer) { +static void print(spirv::VerCapExtAttr triple, DialectAsmPrinter &printer) { auto &os = printer.getStream(); - printer << spirv::TargetEnvAttr::getKindName() << "<" - << spirv::stringifyVersion(targetEnv.getVersion()) << ", ["; - interleaveComma(targetEnv.getExtensionsAttr(), os, [&](Attribute attr) { - os << attr.cast().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().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()) print(targetEnv, printer); + else if (auto vceAttr = attr.dyn_cast()) + 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(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()) - return op->emitError("'") - << symbol - << "' must be a dictionary attribute containing one 32-bit " - "integer attribute 'version', one string array attribute " - "'extensions', one 32-bit integer array attribute " - "'capabilities', and one dictionary attribute 'limits'"; + return op->emitError("'") << symbol << "' must be a spirv::TargetEnvAttr"; } else { return op->emitError("found unsupported '") << symbol << "' attribute on operation"; diff --git a/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp b/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp index f8c5900eb842ae..3743cf44348c57 100644 --- a/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp +++ b/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp @@ -14,95 +14,123 @@ using namespace mlir; +//===----------------------------------------------------------------------===// +// DictionaryDict derived attributes +//===----------------------------------------------------------------------===// + namespace mlir { #include "mlir/Dialect/SPIRV/TargetAndABI.cpp.inc" +//===----------------------------------------------------------------------===// +// Attribute storage classes +//===----------------------------------------------------------------------===// + namespace spirv { namespace detail { +struct VerCapExtAttributeStorage : public AttributeStorage { + using KeyTy = std::tuple; + + 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(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; + using KeyTy = std::pair; - 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(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 extensions, - ArrayRef capabilities, DictionaryAttr limits) { - Builder b(limits.getContext()); +//===----------------------------------------------------------------------===// +// VerCapExtAttr +//===----------------------------------------------------------------------===// - auto versionAttr = b.getI32IntegerAttr(static_cast(version)); +spirv::VerCapExtAttr spirv::VerCapExtAttr::get( + spirv::Version version, ArrayRef capabilities, + ArrayRef extensions, MLIRContext *context) { + Builder b(context); - SmallVector extAttrs; - extAttrs.reserve(extensions.size()); - for (spirv::Extension ext : extensions) - extAttrs.push_back(b.getStringAttr(spirv::stringifyExtension(ext))); + auto versionAttr = b.getI32IntegerAttr(static_cast(version)); SmallVector capAttrs; capAttrs.reserve(capabilities.size()); for (spirv::Capability cap : capabilities) capAttrs.push_back(b.getI32IntegerAttr(static_cast(cap))); - return get(versionAttr, b.getArrayAttr(extAttrs), b.getArrayAttr(capAttrs), - limits); + SmallVector 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( getImpl()->version.cast().getValue().getZExtValue()); } -spirv::TargetEnvAttr::ext_iterator::ext_iterator(ArrayAttr::iterator it) +spirv::VerCapExtAttr::ext_iterator::ext_iterator(ArrayAttr::iterator it) : llvm::mapped_iterator( it, [](Attribute attr) { return *symbolizeExtension(attr.cast().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(); } -spirv::TargetEnvAttr::cap_iterator::cap_iterator(ArrayAttr::iterator it) +spirv::VerCapExtAttr::cap_iterator::cap_iterator(ArrayAttr::iterator it) : llvm::mapped_iterator( it, [](Attribute attr) { @@ -110,25 +138,29 @@ spirv::TargetEnvAttr::cap_iterator::cap_iterator(ArrayAttr::iterator it) attr.cast().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(); } -spirv::ResourceLimitsAttr spirv::TargetEnvAttr::getResourceLimits() { - return getImpl()->limits.cast(); -} - -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()) + 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()) if (spirv::symbolizeExtension(strAttr.getValue())) @@ -137,20 +169,62 @@ LogicalResult spirv::TargetEnvAttr::verifyConstructionInvariants( })) return emitError(loc, "unknown extension in extension list"); - if (!llvm::all_of(capabilities.getValue(), [](Attribute attr) { - if (auto intAttr = attr.dyn_cast()) - 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::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(); +} + +LogicalResult spirv::TargetEnvAttr::verifyConstructionInvariants( + Location loc, spirv::VerCapExtAttr triple, DictionaryAttr limits) { if (!limits.isa()) 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(spirv::Version::V_1_0)), - builder.getI32ArrayAttr({}), - builder.getI32ArrayAttr( - {static_cast(spirv::Capability::Shader)}), - spirv::getDefaultResourceLimits(context)); + auto triple = spirv::VerCapExtAttr::get(spirv::Version::V_1_0, + {spirv::Capability::Shader}, + ArrayRef(), context); + return spirv::TargetEnvAttr::get(triple, + spirv::getDefaultResourceLimits(context)); } spirv::TargetEnvAttr spirv::lookupTargetEnvOrDefault(Operation *op) { diff --git a/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir b/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir index 6caaf8a00a3626..cebd541977ef1c 100644 --- a/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir +++ b/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir @@ -16,7 +16,7 @@ module attributes { spv.target_env = #spv.target_env< - V_1_3, [], [Shader, GroupNonUniformArithmetic], + #spv.vce, { 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, { 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, { 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, { max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32> diff --git a/mlir/test/Dialect/SPIRV/availability.mlir b/mlir/test/Dialect/SPIRV/availability.mlir index 381754c74609e2..a5203a0e4a2ae9 100644 --- a/mlir/test/Dialect/SPIRV/availability.mlir +++ b/mlir/test/Dialect/SPIRV/availability.mlir @@ -2,8 +2,8 @@ // CHECK-LABEL: iadd func @iadd(%arg: i32) -> i32 { - // CHECK: min version: V_1_0 - // CHECK: max version: V_1_5 + // CHECK: min version: v1.0 + // CHECK: max version: v1.5 // CHECK: extensions: [ ] // CHECK: capabilities: [ ] %0 = spv.IAdd %arg, %arg: i32 @@ -12,8 +12,8 @@ func @iadd(%arg: i32) -> i32 { // CHECK: atomic_compare_exchange_weak func @atomic_compare_exchange_weak(%ptr: !spv.ptr, %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 @@ -22,8 +22,8 @@ func @atomic_compare_exchange_weak(%ptr: !spv.ptr, %value: i32, // CHECK-LABEL: subgroup_ballot func @subgroup_ballot(%predicate: i1) -> vector<4xi32> { - // CHECK: min version: V_1_3 - // CHECK: max version: V_1_5 + // CHECK: min version: v1.3 + // CHECK: max version: v1.5 // CHECK: extensions: [ ] // CHECK: capabilities: [ [GroupNonUniformBallot] ] %0 = spv.GroupNonUniformBallot "Workgroup" %predicate : vector<4xi32> @@ -32,8 +32,8 @@ func @subgroup_ballot(%predicate: i1) -> vector<4xi32> { // CHECK-LABEL: module_logical_glsl450 func @module_logical_glsl450() { - // CHECK: spv.module min version: V_1_0 - // CHECK: spv.module max version: V_1_5 + // CHECK: spv.module min version: v1.0 + // CHECK: spv.module max version: v1.5 // CHECK: spv.module extensions: [ ] // CHECK: spv.module capabilities: [ [Shader] ] spv.module "Logical" "GLSL450" { } @@ -42,8 +42,8 @@ func @module_logical_glsl450() { // CHECK-LABEL: module_physical_storage_buffer64_vulkan func @module_physical_storage_buffer64_vulkan() { - // CHECK: spv.module min version: V_1_0 - // CHECK: spv.module max version: V_1_5 + // CHECK: spv.module min version: v1.0 + // CHECK: spv.module max version: v1.5 // CHECK: spv.module extensions: [ [SPV_EXT_physical_storage_buffer, SPV_KHR_physical_storage_buffer] [SPV_KHR_vulkan_memory_model] ] // CHECK: spv.module capabilities: [ [PhysicalStorageBufferAddresses] [VulkanMemoryModel] ] spv.module "PhysicalStorageBuffer64" "Vulkan" { } diff --git a/mlir/test/Dialect/SPIRV/target-and-abi.mlir b/mlir/test/Dialect/SPIRV/target-and-abi.mlir index 1182e28424e4c8..a28ca29e0ab9e0 100644 --- a/mlir/test/Dialect/SPIRV/target-and-abi.mlir +++ b/mlir/test/Dialect/SPIRV/target-and-abi.mlir @@ -106,87 +106,99 @@ func @interface_var() -> (f32 {spv.interface_var_abi = { // spv.target_env //===----------------------------------------------------------------------===// -func @target_env_wrong_type() attributes { - // expected-error @+1 {{expected valid keyword}} - spv.target_env = #spv.target_env<64> +func @target_env_missing_limits() attributes { + spv.target_env = #spv.target_env< + #spv.vce, + // 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 +func @target_env_wrong_limits() attributes { + spv.target_env = #spv.target_env< + #spv.vce, + // 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 +func @target_env() attributes { + // CHECK: spv.target_env = #spv.target_env< + // CHECK-SAME: #spv.vce, + // 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, + { + 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 +func @target_env_extra_fields() attributes { + // expected-error @+6 {{expected '>'}} + spv.target_env = #spv.target_env< + #spv.vce, + { + 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 +//===----------------------------------------------------------------------===// +// 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 +func @vce_missing_fields() attributes { + // expected-error @+1 {{expected ','}} + vce = #spv.vce } { 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 } { 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 } { return } // ----- -func @target_env() attributes { +func @vce_wrong_extension() attributes { + // expected-error @+1 {{unknown extension: SPV_Something}} + vce = #spv.vce +} { return } - // CHECK: spv.target_env = #spv.target_env : 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 } { 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 + vce = #spv.vce } { return } diff --git a/mlir/test/Dialect/SPIRV/target-env.mlir b/mlir/test/Dialect/SPIRV/target-env.mlir index 1e43ec9fbb9d4b..32f36e96f5ea79 100644 --- a/mlir/test/Dialect/SPIRV/target-env.mlir +++ b/mlir/test/Dialect/SPIRV/target-env.mlir @@ -35,7 +35,7 @@ // CHECK-LABEL: @cmp_exchange_weak_suitable_version_capabilities func @cmp_exchange_weak_suitable_version_capabilities(%ptr: !spv.ptr, %value: i32, %comparator: i32) -> i32 attributes { - spv.target_env = #spv.target_env : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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, i32) -> (i32) @@ -44,7 +44,7 @@ func @cmp_exchange_weak_suitable_version_capabilities(%ptr: !spv.ptr, %value: i32, %comparator: i32) -> i32 attributes { - spv.target_env = #spv.target_env : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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, i32) -> (i32) @@ -57,7 +57,7 @@ func @cmp_exchange_weak_unsupported_version(%ptr: !spv.ptr, %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 : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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 : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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, %value: i32, %comparator: i32) -> i32 attributes { - spv.target_env = #spv.target_env : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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, i32) -> (i32) @@ -88,7 +88,7 @@ func @cmp_exchange_weak_missing_capability_kernel(%ptr: !spv.ptr // CHECK-LABEL: @cmp_exchange_weak_missing_capability_atomic_storage func @cmp_exchange_weak_missing_capability_atomic_storage(%ptr: !spv.ptr, %value: i32, %comparator: i32) -> i32 attributes { - spv.target_env = #spv.target_env : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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, i32) -> (i32) @@ -97,7 +97,7 @@ func @cmp_exchange_weak_missing_capability_atomic_storage(%ptr: !spv.ptr vector<4xi32> attributes { - spv.target_env = #spv.target_env : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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 : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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 : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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 : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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 : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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 : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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 : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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 : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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 : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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 : vector<3xi32>}> + spv.target_env = #spv.target_env<#spv.vce, {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"() : () -> ()