diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td @@ -2942,6 +2942,18 @@ [SPV_ISUI_SamplerUnknown, SPV_ISUI_NeedSampler, SPV_ISUI_NoSampler]>; //===----------------------------------------------------------------------===// +// SPIR-V attribute definitions +//===----------------------------------------------------------------------===// + +def SPV_VerCapExtAttr : Attr< + CPred<"$_self.isa<::mlir::spirv::VerCapExtAttr>()">, + "version-capability-extension attribute"> { + let storageType = "::mlir::spirv::VerCapExtAttr"; + let returnType = "::mlir::spirv::VerCapExtAttr"; + let convertFromStorage = "$_self"; +} + +//===----------------------------------------------------------------------===// // SPIR-V type definitions //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVBinaryUtils.h b/mlir/include/mlir/Dialect/SPIRV/SPIRVBinaryUtils.h --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVBinaryUtils.h +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVBinaryUtils.h @@ -34,8 +34,10 @@ #define GET_SPIRV_SERIALIZATION_UTILS #include "mlir/Dialect/SPIRV/SPIRVSerialization.inc" -/// Appends a SPRI-V module header to `header` with the given `idBound`. -void appendModuleHeader(SmallVectorImpl &header, uint32_t idBound); +/// Appends a SPRI-V module header to `header` with the given `version` and +/// `idBound`. +void appendModuleHeader(SmallVectorImpl &header, + spirv::Version version, uint32_t idBound); /// Returns the word-count-prefixed opcode for an SPIR-V instruction. uint32_t getPrefixedOpcode(uint32_t wordCount, spirv::Opcode opcode); diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h b/mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h @@ -23,6 +23,7 @@ class OpBuilder; namespace spirv { +class VerCapExtAttr; // TableGen'erated operation interfaces for querying versions, extensions, and // capabilities. diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td @@ -382,25 +382,25 @@ ### Custom assembly form ``` - addressing-model ::= `"Logical"` | `"Physical32"` | `"Physical64"` - memory-model ::= `"Simple"` | `"GLSL450"` | `"OpenCL"` | `"VulkanKHR"` + addressing-model ::= `Logical` | `Physical32` | `Physical64` | ... + memory-model ::= `Simple` | `GLSL450` | `OpenCL` | `Vulkan` | ... spv-module-op ::= `spv.module` addressing-model memory-model - region + (requires spirv-vce-attribute)? (`attributes` attribute-dict)? + region ``` For example: ``` - spv.module "Logical" "VulkanKHR" { } + spv.module Logical GLSL450 {} - spv.module "Logical" "VulkanKHR" { - func @do_nothing() -> () { + spv.module Logical Vulkan + requires #spv.vce + attributes { some_additional_attr = ... } { + spv.func @do_nothing() -> () { spv.Return } - } attributes { - capability = ["Shader"], - extension = ["SPV_KHR_16bit_storage"] } ``` }]; @@ -408,26 +408,19 @@ let arguments = (ins SPV_AddressingModelAttr:$addressing_model, SPV_MemoryModelAttr:$memory_model, - OptionalAttr:$capabilities, - OptionalAttr:$extensions, - OptionalAttr:$extended_instruction_sets + OptionalAttr:$vce_triple ); let results = (outs); let regions = (region SizedRegion<1>:$body); - let builders = - [OpBuilder<"Builder *, OperationState &state">, - OpBuilder<[{Builder *, OperationState &state, - IntegerAttr addressing_model, - IntegerAttr memory_model}]>, - OpBuilder<[{Builder *, OperationState &state, - spirv::AddressingModel addressing_model, - spirv::MemoryModel memory_model, - /*optional*/ ArrayRef capabilities = {}, - /*optional*/ ArrayRef extensions = {}, - /*optional*/ ArrayAttr extended_instruction_sets = nullptr}]>]; + let builders = [ + OpBuilder<[{Builder *, OperationState &state}]>, + OpBuilder<[{Builder *, OperationState &state, + spirv::AddressingModel addressing_model, + spirv::MemoryModel memory_model}]> + ]; // We need to ensure the block inside the region is properly terminated; // the auto-generated builders do not guarantee that. @@ -438,6 +431,8 @@ let autogenSerialization = 0; let extraClassDeclaration = [{ + static StringRef getVCETripleAttrName() { return "vce_triple"; } + Block& getBlock() { return this->getOperation()->getRegion(0).front(); } diff --git a/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp --- a/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp +++ b/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp @@ -376,13 +376,10 @@ PatternMatchResult GPUModuleConversion::matchAndRewrite( gpu::GPUModuleOp moduleOp, ArrayRef operands, ConversionPatternRewriter &rewriter) const { - // TODO : Generalize this to account for different extensions, - // capabilities, extended_instruction_sets, other addressing models - // and memory models. auto spvModule = rewriter.create( moduleOp.getLoc(), spirv::AddressingModel::Logical, - spirv::MemoryModel::GLSL450, spirv::Capability::Shader, - spirv::Extension::SPV_KHR_storage_buffer_storage_class); + spirv::MemoryModel::GLSL450); + // Move the region from the module op into the SPIR-V module. Region &spvModuleRegion = spvModule.body(); rewriter.inlineRegionBefore(moduleOp.body(), spvModuleRegion, diff --git a/mlir/lib/Dialect/SPIRV/SPIRVOps.cpp b/mlir/lib/Dialect/SPIRV/SPIRVOps.cpp --- a/mlir/lib/Dialect/SPIRV/SPIRVOps.cpp +++ b/mlir/lib/Dialect/SPIRV/SPIRVOps.cpp @@ -12,6 +12,7 @@ #include "mlir/Dialect/SPIRV/SPIRVOps.h" +#include "mlir/Dialect/SPIRV/SPIRVAttributes.h" #include "mlir/Dialect/SPIRV/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/SPIRVTypes.h" #include "mlir/IR/Builders.h" @@ -97,10 +98,12 @@ return builder.getStrArrayAttr(enumValStrs); } +/// Parses the next string attribute in `parser` as an enumerant of the given +/// `EnumClass`. template static ParseResult -parseEnumAttribute(EnumClass &value, OpAsmParser &parser, - StringRef attrName = spirv::attributeName()) { +parseEnumStrAttr(EnumClass &value, OpAsmParser &parser, + StringRef attrName = spirv::attributeName()) { Attribute attrVal; SmallVector attr; auto loc = parser.getCurrentLocation(); @@ -122,11 +125,49 @@ return success(); } +/// Parses the next string attribute in `parser` as an enumerant of the given +/// `EnumClass` and inserts the enumerant into `state` as an 32-bit integer +/// attribute with the enum class's name as attribute name. template static ParseResult -parseEnumAttribute(EnumClass &value, OpAsmParser &parser, OperationState &state, - StringRef attrName = spirv::attributeName()) { - if (parseEnumAttribute(value, parser)) { +parseEnumStrAttr(EnumClass &value, OpAsmParser &parser, OperationState &state, + StringRef attrName = spirv::attributeName()) { + if (parseEnumStrAttr(value, parser)) { + return failure(); + } + state.addAttribute(attrName, parser.getBuilder().getI32IntegerAttr( + llvm::bit_cast(value))); + return success(); +} + +/// Parses the next keyword in `parser` as an enumerant of the given +/// `EnumClass`. +template +static ParseResult +parseEnumKeywordAttr(EnumClass &value, OpAsmParser &parser, + StringRef attrName = spirv::attributeName()) { + StringRef keyword; + SmallVector attr; + auto loc = parser.getCurrentLocation(); + if (parser.parseKeyword(&keyword)) + return failure(); + if (Optional attr = spirv::symbolizeEnum()(keyword)) { + value = attr.getValue(); + return success(); + } + return parser.emitError(loc, "invalid ") + << attrName << " attribute specification: " << keyword; +} + +/// Parses the next keyword in `parser` as an enumerant of the given `EnumClass` +/// and inserts the enumerant into `state` as an 32-bit integer attribute with +/// the enum class's name as attribute name. +template +static ParseResult +parseEnumKeywordAttr(EnumClass &value, OpAsmParser &parser, + OperationState &state, + StringRef attrName = spirv::attributeName()) { + if (parseEnumKeywordAttr(value, parser)) { return failure(); } state.addAttribute(attrName, parser.getBuilder().getI32IntegerAttr( @@ -143,7 +184,7 @@ } spirv::MemoryAccess memoryAccessAttr; - if (parseEnumAttribute(memoryAccessAttr, parser, state)) { + if (parseEnumStrAttr(memoryAccessAttr, parser, state)) { return failure(); } @@ -463,8 +504,8 @@ OpAsmParser::OperandType ptrInfo, valueInfo; Type type; llvm::SMLoc loc; - if (parseEnumAttribute(scope, parser, state, kMemoryScopeAttrName) || - parseEnumAttribute(memoryScope, parser, state, kSemanticsAttrName) || + if (parseEnumStrAttr(scope, parser, state, kMemoryScopeAttrName) || + parseEnumStrAttr(memoryScope, parser, state, kSemanticsAttrName) || parser.parseOperandList(operandInfo, (hasValue ? 2 : 1)) || parser.getCurrentLocation(&loc) || parser.parseColonType(type)) return failure(); @@ -521,10 +562,10 @@ spirv::Scope executionScope; spirv::GroupOperation groupOperation; OpAsmParser::OperandType valueInfo; - if (parseEnumAttribute(executionScope, parser, state, - kExecutionScopeAttrName) || - parseEnumAttribute(groupOperation, parser, state, - kGroupOperationAttrName) || + if (parseEnumStrAttr(executionScope, parser, state, + kExecutionScopeAttrName) || + parseEnumStrAttr(groupOperation, parser, state, + kGroupOperationAttrName) || parser.parseOperand(valueInfo)) return failure(); @@ -845,11 +886,11 @@ spirv::MemorySemantics equalSemantics, unequalSemantics; SmallVector operandInfo; Type type; - if (parseEnumAttribute(memoryScope, parser, state, kMemoryScopeAttrName) || - parseEnumAttribute(equalSemantics, parser, state, - kEqualSemanticsAttrName) || - parseEnumAttribute(unequalSemantics, parser, state, - kUnequalSemanticsAttrName) || + if (parseEnumStrAttr(memoryScope, parser, state, kMemoryScopeAttrName) || + parseEnumStrAttr(equalSemantics, parser, state, + kEqualSemanticsAttrName) || + parseEnumStrAttr(unequalSemantics, parser, state, + kUnequalSemanticsAttrName) || parser.parseOperandList(operandInfo, 3)) return failure(); @@ -1394,7 +1435,7 @@ SmallVector interfaceVars; FlatSymbolRefAttr fn; - if (parseEnumAttribute(execModel, parser, state) || + if (parseEnumStrAttr(execModel, parser, state) || parser.parseAttribute(fn, Type(), kFnNameAttrName, state.attributes)) { return failure(); } @@ -1452,7 +1493,7 @@ spirv::ExecutionMode execMode; Attribute fn; if (parser.parseAttribute(fn, kFnNameAttrName, state.attributes) || - parseEnumAttribute(execMode, parser, state)) { + parseEnumStrAttr(execMode, parser, state)) { return failure(); } @@ -1515,7 +1556,7 @@ // Parse the optional function control keyword. spirv::FunctionControl fnControl; - if (parseEnumAttribute(fnControl, parser, state)) + if (parseEnumStrAttr(fnControl, parser, state)) return failure(); // If additional attributes are present, parse them. @@ -1840,8 +1881,7 @@ spirv::StorageClass storageClass; OpAsmParser::OperandType ptrInfo; Type elementType; - if (parseEnumAttribute(storageClass, parser) || - parser.parseOperand(ptrInfo) || + if (parseEnumStrAttr(storageClass, parser) || parser.parseOperand(ptrInfo) || parseMemoryAccessAttributes(parser, state) || parser.parseOptionalAttrDict(state.attributes) || parser.parseColon() || parser.parseType(elementType)) { @@ -2068,38 +2108,15 @@ ensureTerminator(*state.addRegion(), *builder, state.location); } -// TODO(ravishankarm): This is only here for resolving some dependency outside -// of mlir. Remove once it is done. -void spirv::ModuleOp::build(Builder *builder, OperationState &state, - IntegerAttr addressing_model, - IntegerAttr memory_model) { - state.addAttribute("addressing_model", addressing_model); - state.addAttribute("memory_model", memory_model); - build(builder, state); -} - void spirv::ModuleOp::build(Builder *builder, OperationState &state, spirv::AddressingModel addressing_model, - spirv::MemoryModel memory_model, - ArrayRef capabilities, - ArrayRef extensions, - ArrayAttr extended_instruction_sets) { + spirv::MemoryModel memory_model) { state.addAttribute( "addressing_model", builder->getI32IntegerAttr(static_cast(addressing_model))); state.addAttribute("memory_model", builder->getI32IntegerAttr( static_cast(memory_model))); - if (!capabilities.empty()) - state.addAttribute("capabilities", - getStrArrayAttrForEnumList( - *builder, capabilities, spirv::stringifyCapability)); - if (!extensions.empty()) - state.addAttribute("extensions", - getStrArrayAttrForEnumList( - *builder, extensions, spirv::stringifyExtension)); - if (extended_instruction_sets) - state.addAttribute("extended_instruction_sets", extended_instruction_sets); - build(builder, state); + ensureTerminator(*state.addRegion(), *builder, state.location); } static ParseResult parseModuleOp(OpAsmParser &parser, OperationState &state) { @@ -2108,15 +2125,22 @@ // Parse attributes spirv::AddressingModel addrModel; spirv::MemoryModel memoryModel; - if (parseEnumAttribute(addrModel, parser, state) || - parseEnumAttribute(memoryModel, parser, state)) { + if (parseEnumKeywordAttr(addrModel, parser, state) || + parseEnumKeywordAttr(memoryModel, parser, state)) return failure(); + + if (succeeded(parser.parseOptionalKeyword("requires"))) { + spirv::VerCapExtAttr vceTriple; + if (parser.parseAttribute(vceTriple, + spirv::ModuleOp::getVCETripleAttrName(), + state.attributes)) + return failure(); } - if (parser.parseRegion(*body, /*arguments=*/{}, /*argTypes=*/{})) + if (parser.parseOptionalAttrDictWithKeyword(state.attributes)) return failure(); - if (parser.parseOptionalAttrDictWithKeyword(state.attributes)) + if (parser.parseRegion(*body, /*arguments=*/{}, /*argTypes=*/{})) return failure(); spirv::ModuleOp::ensureTerminator(*body, parser.getBuilder(), state.location); @@ -2126,35 +2150,32 @@ static void print(spirv::ModuleOp moduleOp, OpAsmPrinter &printer) { printer << spirv::ModuleOp::getOperationName(); - // Only print out addressing model and memory model in a nicer way if both - // presents. Otherwise, print them in the general form. This helps - // debugging ill-formed ModuleOp. SmallVector elidedAttrs; + + printer << " " << spirv::stringifyAddressingModel(moduleOp.addressing_model()) + << " " << spirv::stringifyMemoryModel(moduleOp.memory_model()); auto addressingModelAttrName = spirv::attributeName(); auto memoryModelAttrName = spirv::attributeName(); - if (moduleOp.getAttr(addressingModelAttrName) && - moduleOp.getAttr(memoryModelAttrName)) { - printer << " \"" - << spirv::stringifyAddressingModel(moduleOp.addressing_model()) - << "\" \"" << spirv::stringifyMemoryModel(moduleOp.memory_model()) - << '"'; - elidedAttrs.assign({addressingModelAttrName, memoryModelAttrName}); + elidedAttrs.assign({addressingModelAttrName, memoryModelAttrName}); + + if (Optional triple = moduleOp.vce_triple()) { + printer << " requires " << *triple; + elidedAttrs.push_back(spirv::ModuleOp::getVCETripleAttrName()); } + printer.printOptionalAttrDictWithKeyword(moduleOp.getAttrs(), elidedAttrs); printer.printRegion(moduleOp.body(), /*printEntryBlockArgs=*/false, /*printBlockTerminators=*/false); - printer.printOptionalAttrDictWithKeyword(moduleOp.getAttrs(), elidedAttrs); } static LogicalResult verify(spirv::ModuleOp moduleOp) { auto &op = *moduleOp.getOperation(); auto *dialect = op.getDialect(); - auto &body = op.getRegion(0).front(); DenseMap, spirv::EntryPointOp> entryPoints; SymbolTable table(moduleOp); - for (auto &op : body) { + for (auto &op : moduleOp.getBlock()) { if (op.getDialect() != dialect) return op.emitError("'spv.module' can only contain spv.* ops"); @@ -2207,26 +2228,6 @@ } } - // Verify capabilities. ODS already guarantees that we have an array of - // string attributes. - if (auto caps = moduleOp.getAttrOfType("capabilities")) { - for (auto cap : caps.getValue()) { - auto capStr = cap.cast().getValue(); - if (!spirv::symbolizeCapability(capStr)) - return moduleOp.emitOpError("uses unknown capability: ") << capStr; - } - } - - // Verify extensions. ODS already guarantees that we have an array of - // string attributes. - if (auto exts = moduleOp.getAttrOfType("extensions")) { - for (auto ext : exts.getValue()) { - auto extStr = ext.cast().getValue(); - if (!spirv::symbolizeExtension(extStr)) - return moduleOp.emitOpError("uses unknown extension: ") << extStr; - } - } - return success(); } @@ -2479,7 +2480,7 @@ SmallVector operandInfo; auto loc = parser.getCurrentLocation(); Type elementType; - if (parseEnumAttribute(storageClass, parser) || + if (parseEnumStrAttr(storageClass, parser) || parser.parseOperandList(operandInfo, 2) || parseMemoryAccessAttributes(parser, state) || parser.parseColon() || parser.parseType(elementType)) { diff --git a/mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp b/mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp --- a/mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp +++ b/mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp @@ -12,6 +12,7 @@ #include "mlir/Dialect/SPIRV/Serialization.h" +#include "mlir/Dialect/SPIRV/SPIRVAttributes.h" #include "mlir/Dialect/SPIRV/SPIRVBinaryUtils.h" #include "mlir/Dialect/SPIRV/SPIRVOps.h" #include "mlir/Dialect/SPIRV/SPIRVTypes.h" @@ -106,9 +107,6 @@ /// in the deserializer. LogicalResult processCapability(ArrayRef operands); - /// Attaches all collected capabilities to `module` as an attribute. - void attachCapabilities(); - /// Processes the SPIR-V OpExtension with `operands` and updates bookkeeping /// in the deserializer. LogicalResult processExtension(ArrayRef words); @@ -117,8 +115,9 @@ /// bookkeeping in the deserializer. LogicalResult processExtInstImport(ArrayRef words); - /// Attaches all collected extensions to `module` as an attribute. - void attachExtensions(); + /// Attaches (version, capabilities, extensions) triple to `module` as an + /// attribute. + void attachVCETriple(); /// Processes the SPIR-V OpMemoryModel with `operands` and updates `module`. LogicalResult processMemoryModel(ArrayRef operands); @@ -397,11 +396,13 @@ OpBuilder opBuilder; + spirv::Version version; + /// The list of capabilities used by the module. llvm::SmallSetVector capabilities; /// The list of extensions used by the module. - llvm::SmallSetVector extensions; + llvm::SmallSetVector extensions; // Result to type mapping. DenseMap typeMap; @@ -507,9 +508,7 @@ } } - // Attaches the capabilities/extensions as an attribute to the module. - attachCapabilities(); - attachExtensions(); + attachVCETriple(); LLVM_DEBUG(llvm::dbgs() << "+++ completed deserialization +++\n"); return success(); @@ -524,9 +523,6 @@ spirv::ModuleOp Deserializer::createModuleOp() { Builder builder(context); OperationState state(unknownLoc, spirv::ModuleOp::getOperationName()); - // TODO(antiagainst): use target environment to select the version - state.addAttribute("major_version", builder.getI32IntegerAttr(1)); - state.addAttribute("minor_version", builder.getI32IntegerAttr(0)); spirv::ModuleOp::build(&builder, state); return cast(Operation::create(state)); } @@ -539,6 +535,32 @@ if (binary[0] != spirv::kMagicNumber) return emitError(unknownLoc, "incorrect magic number"); + // Version number bytes: 0 | major number | minor number | 0 + uint32_t majorVersion = (binary[1] << 8) >> 24; + uint32_t minorVersion = (binary[1] << 16) >> 24; + if (majorVersion == 1) { + switch (minorVersion) { +#define MIN_VERSION_CASE(v) \ + case v: \ + version = spirv::Version::V_1_##v; \ + break + + MIN_VERSION_CASE(0); + MIN_VERSION_CASE(1); + MIN_VERSION_CASE(2); + MIN_VERSION_CASE(3); + MIN_VERSION_CASE(4); + MIN_VERSION_CASE(5); +#undef MIN_VERSION_CASE + default: + return emitError(unknownLoc, "unspported SPIR-V minor version: ") + << minorVersion; + } + } else { + return emitError(unknownLoc, "unspported SPIR-V major version: ") + << majorVersion; + } + // TODO(antiagainst): generator number, bound, schema curOffset = spirv::kHeaderWordCount; return success(); @@ -556,20 +578,6 @@ return success(); } -void Deserializer::attachCapabilities() { - if (capabilities.empty()) - return; - - SmallVector caps; - caps.reserve(capabilities.size()); - - for (auto cap : capabilities) { - caps.push_back(spirv::stringifyCapability(cap)); - } - - module->setAttr("capabilities", opBuilder.getStrArrayAttr(caps)); -} - LogicalResult Deserializer::processExtension(ArrayRef words) { if (words.empty()) { return emitError( @@ -579,12 +587,14 @@ unsigned wordIndex = 0; StringRef extName = decodeStringLiteral(words, wordIndex); - if (wordIndex != words.size()) { + if (wordIndex != words.size()) return emitError(unknownLoc, "unexpected trailing words in OpExtension instruction"); - } + auto ext = spirv::symbolizeExtension(extName); + if (!ext) + return emitError(unknownLoc, "unknown extension: ") << extName; - extensions.insert(extName); + extensions.insert(*ext); return success(); } @@ -604,12 +614,10 @@ return success(); } -void Deserializer::attachExtensions() { - if (extensions.empty()) - return; - - module->setAttr("extensions", - opBuilder.getStrArrayAttr(extensions.getArrayRef())); +void Deserializer::attachVCETriple() { + module->setAttr(spirv::ModuleOp::getVCETripleAttrName(), + spirv::VerCapExtAttr::get(version, capabilities.getArrayRef(), + extensions.getArrayRef(), context)); } LogicalResult Deserializer::processMemoryModel(ArrayRef operands) { diff --git a/mlir/lib/Dialect/SPIRV/Serialization/SPIRVBinaryUtils.cpp b/mlir/lib/Dialect/SPIRV/Serialization/SPIRVBinaryUtils.cpp --- a/mlir/lib/Dialect/SPIRV/Serialization/SPIRVBinaryUtils.cpp +++ b/mlir/lib/Dialect/SPIRV/Serialization/SPIRVBinaryUtils.cpp @@ -11,15 +11,28 @@ //===----------------------------------------------------------------------===// #include "mlir/Dialect/SPIRV/SPIRVBinaryUtils.h" +#include "mlir/Dialect/SPIRV/SPIRVTypes.h" using namespace mlir; void spirv::appendModuleHeader(SmallVectorImpl &header, - uint32_t idBound) { - // The major and minor version number for the generated SPIR-V binary. - // TODO(antiagainst): use target environment to select the version - constexpr uint8_t kMajorVersion = 1; - constexpr uint8_t kMinorVersion = 0; + spirv::Version version, uint32_t idBound) { + uint32_t majorVersion = 1; + uint32_t minorVersion = 0; + switch (version) { +#define MIN_VERSION_CASE(v) \ + case spirv::Version::V_1_##v: \ + minorVersion = v; \ + break + + MIN_VERSION_CASE(0); + MIN_VERSION_CASE(1); + MIN_VERSION_CASE(2); + MIN_VERSION_CASE(3); + MIN_VERSION_CASE(4); + MIN_VERSION_CASE(5); +#undef MIN_VERSION_CASE + } // See "2.3. Physical Layout of a SPIR-V Module and Instruction" in the SPIR-V // spec for the definition of the binary module header. @@ -37,7 +50,7 @@ // | 0 (reserved for instruction schema) | // +-------------------------------------------------------------------------+ header.push_back(spirv::kMagicNumber); - header.push_back((kMajorVersion << 16) | (kMinorVersion << 8)); + header.push_back((majorVersion << 16) | (minorVersion << 8)); header.push_back(kGeneratorNumber); header.push_back(idBound); // bound header.push_back(0); // Schema (reserved word) diff --git a/mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp b/mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp --- a/mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp +++ b/mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp @@ -13,6 +13,7 @@ #include "mlir/Dialect/SPIRV/Serialization.h" #include "mlir/ADT/TypeSwitch.h" +#include "mlir/Dialect/SPIRV/SPIRVAttributes.h" #include "mlir/Dialect/SPIRV/SPIRVBinaryUtils.h" #include "mlir/Dialect/SPIRV/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/SPIRVOps.h" @@ -490,7 +491,7 @@ binary.clear(); binary.reserve(moduleSize); - spirv::appendModuleHeader(binary, nextID); + spirv::appendModuleHeader(binary, module.vce_triple()->getVersion(), nextID); binary.append(capabilities.begin(), capabilities.end()); binary.append(extensions.begin(), extensions.end()); binary.append(extendedSets.begin(), extendedSets.end()); @@ -536,28 +537,16 @@ } void Serializer::processCapability() { - auto caps = module.getAttrOfType("capabilities"); - if (!caps) - return; - - for (auto cap : caps.getValue()) { - auto capStr = cap.cast().getValue(); - auto capVal = spirv::symbolizeCapability(capStr); + for (auto cap : module.vce_triple()->getCapabilities()) encodeInstructionInto(capabilities, spirv::Opcode::OpCapability, - {static_cast(*capVal)}); - } + {static_cast(cap)}); } void Serializer::processExtension() { - auto exts = module.getAttrOfType("extensions"); - if (!exts) - return; - - SmallVector extName; - for (auto ext : exts.getValue()) { - auto extStr = ext.cast().getValue(); + llvm::SmallVector extName; + for (spirv::Extension ext : module.vce_triple()->getExtensions()) { extName.clear(); - spirv::encodeStringLiteralInto(extName, extStr); + spirv::encodeStringLiteralInto(extName, spirv::stringifyExtension(ext)); encodeInstructionInto(extensions, spirv::Opcode::OpExtension, extName); } } @@ -1812,6 +1801,10 @@ LogicalResult spirv::serialize(spirv::ModuleOp module, SmallVectorImpl &binary) { + if (!module.vce_triple().hasValue()) + return module.emitError( + "module must have 'vce_triple' attribute to be serializeable"); + Serializer serializer(module); if (failed(serializer.serialize())) diff --git a/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp --- a/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp @@ -150,7 +150,7 @@ auto triple = spirv::VerCapExtAttr::get( deducedVersion, deducedCapabilities.getArrayRef(), deducedExtensions.getArrayRef(), &getContext()); - module.setAttr("vce_triple", triple); + module.setAttr(spirv::ModuleOp::getVCETripleAttrName(), triple); } std::unique_ptr> diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir --- a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir @@ -7,7 +7,7 @@ return } - // CHECK-LABEL: spv.module "Logical" "GLSL450" + // CHECK-LABEL: spv.module Logical GLSL450 // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") gpu.module @kernels { gpu.func @builtin_workgroup_id_x() @@ -30,7 +30,7 @@ return } - // CHECK-LABEL: spv.module "Logical" "GLSL450" + // CHECK-LABEL: spv.module Logical GLSL450 // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") gpu.module @kernels { gpu.func @builtin_workgroup_id_y() @@ -53,7 +53,7 @@ return } - // CHECK-LABEL: spv.module "Logical" "GLSL450" + // CHECK-LABEL: spv.module Logical GLSL450 // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") gpu.module @kernels { gpu.func @builtin_workgroup_id_z() @@ -76,7 +76,7 @@ return } - // CHECK-LABEL: spv.module "Logical" "GLSL450" + // CHECK-LABEL: spv.module Logical GLSL450 gpu.module @kernels { gpu.func @builtin_workgroup_size_x() attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[32, 1, 1]>: vector<3xi32>}} { @@ -100,7 +100,7 @@ return } - // CHECK-LABEL: spv.module "Logical" "GLSL450" + // CHECK-LABEL: spv.module Logical GLSL450 gpu.module @kernels { gpu.func @builtin_workgroup_size_y() attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { @@ -121,7 +121,7 @@ return } - // CHECK-LABEL: spv.module "Logical" "GLSL450" + // CHECK-LABEL: spv.module Logical GLSL450 gpu.module @kernels { gpu.func @builtin_workgroup_size_z() attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { @@ -142,7 +142,7 @@ return } - // CHECK-LABEL: spv.module "Logical" "GLSL450" + // CHECK-LABEL: spv.module Logical GLSL450 // CHECK: spv.globalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId") gpu.module @kernels { gpu.func @builtin_local_id_x() @@ -165,7 +165,7 @@ return } - // CHECK-LABEL: spv.module "Logical" "GLSL450" + // CHECK-LABEL: spv.module Logical GLSL450 // CHECK: spv.globalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") gpu.module @kernels { gpu.func @builtin_num_workgroups_x() diff --git a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir --- a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir @@ -15,7 +15,7 @@ return } - // CHECK-LABEL: spv.module "Logical" "GLSL450" + // CHECK-LABEL: spv.module Logical GLSL450 gpu.module @kernels { // CHECK-DAG: spv.globalVariable [[NUMWORKGROUPSVAR:@.*]] built_in("NumWorkgroups") : !spv.ptr, Input> // CHECK-DAG: spv.globalVariable [[LOCALINVOCATIONIDVAR:@.*]] built_in("LocalInvocationId") : !spv.ptr, Input> diff --git a/mlir/test/Conversion/GPUToSPIRV/simple.mlir b/mlir/test/Conversion/GPUToSPIRV/simple.mlir --- a/mlir/test/Conversion/GPUToSPIRV/simple.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/simple.mlir @@ -2,7 +2,7 @@ module attributes {gpu.container_module} { gpu.module @kernels { - // CHECK: spv.module "Logical" "GLSL450" { + // CHECK: spv.module Logical GLSL450 { // CHECK-LABEL: spv.func @basic_module_structure // CHECK-SAME: {{%.*}}: f32 {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}} // CHECK-SAME: {{%.*}}: !spv.ptr [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}} @@ -12,7 +12,6 @@ // CHECK: spv.Return gpu.return } - // CHECK: attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]} } func @main() { diff --git a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir --- a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir +++ b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir @@ -5,7 +5,7 @@ // CHECK: call @vulkanLaunch(%[[index]], %[[index]], %[[index]], %[[index]], %[[index]], %[[index]], %[[resource]]) {spirv_blob = "{{.*}}", spirv_entry_point = "kernel"} module attributes {gpu.container_module} { - spv.module "Logical" "GLSL450" { + spv.module Logical GLSL450 requires #spv.vce { spv.globalVariable @kernel_arg_0 bind(0, 0) : !spv.ptr [0]>, StorageBuffer> spv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} { %0 = spv._address_of @kernel_arg_0 : !spv.ptr [0]>, StorageBuffer> @@ -17,7 +17,7 @@ } spv.EntryPoint "GLCompute" @kernel spv.ExecutionMode @kernel "LocalSize", 1, 1, 1 - } attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]} + } gpu.module @kernels { gpu.func @kernel(%arg0: memref<12xf32>) kernel { gpu.return diff --git a/mlir/test/Dialect/SPIRV/Serialization/arithmetic-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/arithmetic-ops.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/arithmetic-ops.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/arithmetic-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @fmul(%arg0 : f32, %arg1 : f32) "None" { // CHECK: {{%.*}}= spv.FMul {{%.*}}, {{%.*}} : f32 %0 = spv.FMul %arg0, %arg1 : f32 diff --git a/mlir/test/Dialect/SPIRV/Serialization/array.mlir b/mlir/test/Dialect/SPIRV/Serialization/array.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/array.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/array.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @array_stride(%arg0 : !spv.ptr [128]>, StorageBuffer>, %arg1 : i32, %arg2 : i32) "None" { // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr [128]>, StorageBuffer> %2 = spv.AccessChain %arg0[%arg1, %arg2] : !spv.ptr [128]>, StorageBuffer> @@ -10,7 +10,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK: spv.globalVariable {{@.*}} : !spv.ptr, StorageBuffer> spv.globalVariable @var0 : !spv.ptr, StorageBuffer> // CHECK: spv.globalVariable {{@.*}} : !spv.ptr>, Input> diff --git a/mlir/test/Dialect/SPIRV/Serialization/atomic-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/atomic-ops.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/atomic-ops.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/atomic-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @atomic_compare_exchange_weak spv.func @atomic_compare_exchange_weak(%ptr: !spv.ptr, %value: i32, %comparator: i32) -> i32 "None" { // CHECK: spv.AtomicCompareExchangeWeak "Workgroup" "Release" "Acquire" %{{.*}}, %{{.*}}, %{{.*}} : !spv.ptr diff --git a/mlir/test/Dialect/SPIRV/Serialization/barrier.mlir b/mlir/test/Dialect/SPIRV/Serialization/barrier.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/barrier.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/barrier.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @memory_barrier_0() -> () "None" { // CHECK: spv.MemoryBarrier "Device", "Release|UniformMemory" spv.MemoryBarrier "Device", "Release|UniformMemory" diff --git a/mlir/test/Dialect/SPIRV/Serialization/bit-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/bit-ops.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/bit-ops.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/bit-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @bitcount(%arg: i32) -> i32 "None" { // CHECK: spv.BitCount {{%.*}} : i32 %0 = spv.BitCount %arg : i32 diff --git a/mlir/test/Dialect/SPIRV/Serialization/cast-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/cast-ops.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/cast-ops.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/cast-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @bit_cast(%arg0 : f32) "None" { // CHECK: {{%.*}} = spv.Bitcast {{%.*}} : f32 to i32 %0 = spv.Bitcast %arg0 : f32 to i32 @@ -14,7 +14,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @convert_f_to_s(%arg0 : f32) -> i32 "None" { // CHECK: {{%.*}} = spv.ConvertFToS {{%.*}} : f32 to i32 %0 = spv.ConvertFToS %arg0 : f32 to i32 diff --git a/mlir/test/Dialect/SPIRV/Serialization/composite-op.mlir b/mlir/test/Dialect/SPIRV/Serialization/composite-op.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/composite-op.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/composite-op.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @composite_insert(%arg0 : !spv.struct, f32>>, %arg1: !spv.array<4xf32>) -> !spv.struct, f32>> "None" { // CHECK: spv.CompositeInsert {{%.*}}, {{%.*}}[1 : i32, 0 : i32] : !spv.array<4 x f32> into !spv.struct, f32>> %0 = spv.CompositeInsert %arg1, %arg0[1 : i32, 0 : i32] : !spv.array<4xf32> into !spv.struct, f32>> diff --git a/mlir/test/Dialect/SPIRV/Serialization/constant.mlir b/mlir/test/Dialect/SPIRV/Serialization/constant.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/constant.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/constant.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @bool_const spv.func @bool_const() -> () "None" { // CHECK: spv.constant true diff --git a/mlir/test/Dialect/SPIRV/Serialization/entry-point.mlir b/mlir/test/Dialect/SPIRV/Serialization/entry-point.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/entry-point.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/entry-point.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @noop() -> () "None" { spv.Return } @@ -12,7 +12,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK: spv.globalVariable @var2 : !spv.ptr // CHECK-NEXT: spv.globalVariable @var3 : !spv.ptr // CHECK-NEXT: spv.func @noop({{%.*}}: !spv.ptr, {{%.*}}: !spv.ptr) "None" diff --git a/mlir/test/Dialect/SPIRV/Serialization/execution-mode.mlir b/mlir/test/Dialect/SPIRV/Serialization/execution-mode.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/execution-mode.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/execution-mode.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { spv.Return } diff --git a/mlir/test/Dialect/SPIRV/Serialization/function-call.mlir b/mlir/test/Dialect/SPIRV/Serialization/function-call.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/function-call.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/function-call.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.globalVariable @var1 : !spv.ptr, Input> spv.func @fmain() -> i32 "None" { %0 = spv.constant 16 : i32 diff --git a/mlir/test/Dialect/SPIRV/Serialization/global-variable.mlir b/mlir/test/Dialect/SPIRV/Serialization/global-variable.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/global-variable.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/global-variable.mlir @@ -5,7 +5,7 @@ // CHECK-NEXT: spv.globalVariable @var2 built_in("GlobalInvocationId") : !spv.ptr, Input> // CHECK-NEXT: spv.globalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr, Input> -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.globalVariable @var0 bind(1, 0) : !spv.ptr spv.globalVariable @var1 bind(0, 1) : !spv.ptr spv.globalVariable @var2 {built_in = "GlobalInvocationId"} : !spv.ptr, Input> @@ -14,7 +14,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK: spv.globalVariable @var1 : !spv.ptr // CHECK-NEXT: spv.globalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr spv.globalVariable @var1 : !spv.ptr @@ -23,7 +23,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr, Input> spv.func @foo() "None" { // CHECK: %[[ADDR:.*]] = spv._address_of @globalInvocationID : !spv.ptr, Input> diff --git a/mlir/test/Dialect/SPIRV/Serialization/glsl-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/glsl-ops.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/glsl-ops.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/glsl-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @fmul(%arg0 : f32, %arg1 : f32) "None" { // CHECK: {{%.*}} = spv.GLSL.Exp {{%.*}} : f32 %0 = spv.GLSL.Exp %arg0 : f32 diff --git a/mlir/test/Dialect/SPIRV/Serialization/group-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/group-ops.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/group-ops.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/group-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @subgroup_ballot spv.func @subgroup_ballot(%predicate: i1) -> vector<4xi32> "None" { // CHECK: %{{.*}} = spv.SubgroupBallotKHR %{{.*}}: vector<4xi32> diff --git a/mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @iequal_scalar(%arg0: i32, %arg1: i32) "None" { // CHECK: {{.*}} = spv.IEqual {{.*}}, {{.*}} : i32 %0 = spv.IEqual %arg0, %arg1 : i32 @@ -82,7 +82,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.specConstant @condition_scalar = true spv.func @select() -> () "None" { %0 = spv.constant 4.0 : f32 diff --git a/mlir/test/Dialect/SPIRV/Serialization/loop.mlir b/mlir/test/Dialect/SPIRV/Serialization/loop.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/loop.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/loop.mlir @@ -2,7 +2,7 @@ // Single loop -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // for (int i = 0; i < count; ++i) {} spv.func @loop(%count : i32) -> () "None" { %zero = spv.constant 0: i32 @@ -55,13 +55,11 @@ spv.Return } spv.EntryPoint "GLCompute" @main -} attributes { - capabilities = ["Shader"] } // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.globalVariable @GV1 bind(0, 0) : !spv.ptr [0]>, StorageBuffer> spv.globalVariable @GV2 bind(0, 1) : !spv.ptr [0]>, StorageBuffer> spv.func @loop_kernel() "None" { @@ -103,13 +101,13 @@ } spv.EntryPoint "GLCompute" @loop_kernel spv.ExecutionMode @loop_kernel "LocalSize", 1, 1, 1 -} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]} +} // ----- // Nested loop -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // for (int i = 0; i < count; ++i) { // for (int j = 0; j < count; ++j) { } // } @@ -207,7 +205,5 @@ spv.Return } spv.EntryPoint "GLCompute" @main -} attributes { - capabilities = ["Shader"] } diff --git a/mlir/test/Dialect/SPIRV/Serialization/memory-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/memory-ops.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/memory-ops.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/memory-ops.mlir @@ -4,7 +4,7 @@ // CHECK-NEXT: [[VALUE:%.*]] = spv.Load "Input" [[ARG1]] : f32 // CHECK-NEXT: spv.Store "Output" [[ARG2]], [[VALUE]] : f32 -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @load_store(%arg0 : !spv.ptr, %arg1 : !spv.ptr) "None" { %1 = spv.Load "Input" %arg0 : f32 spv.Store "Output" %arg1, %1 : f32 @@ -14,7 +14,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @access_chain(%arg0 : !spv.ptr>, Function>, %arg1 : i32, %arg2 : i32) "None" { // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr>, Function> // CHECK-NEXT: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr>, Function> @@ -26,7 +26,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @load_store_zero_rank_float(%arg0: !spv.ptr [0]>, StorageBuffer>, %arg1: !spv.ptr [0]>, StorageBuffer>) "None" { // CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr [0]> // CHECK-NEXT: [[VAL:%.*]] = spv.Load "StorageBuffer" [[LOAD_PTR]] : f32 diff --git a/mlir/test/Dialect/SPIRV/Serialization/module.mlir b/mlir/test/Dialect/SPIRV/Serialization/module.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/module.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/module.mlir @@ -1,12 +1,12 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -// CHECK: spv.module "Logical" "GLSL450" { +// CHECK: spv.module Logical GLSL450 requires #spv.vce { // CHECK-NEXT: spv.func @foo() "None" { // CHECK-NEXT: spv.Return // CHECK-NEXT: } -// CHECK-NEXT: } attributes {major_version = 1 : i32, minor_version = 0 : i32} +// CHECK-NEXT: } -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { spv.Return } @@ -14,17 +14,19 @@ // ----- -spv.module "Logical" "GLSL450" { -} attributes { - // CHECK: capabilities = ["Shader", "Float16"] - capabilities = ["Shader", "Float16"] +// CHECK: v1.5 +spv.module Logical GLSL450 requires #spv.vce { } // ----- -spv.module "Logical" "GLSL450" { -} attributes { - // CHECK: extensions = ["SPV_KHR_float_controls", "SPV_KHR_subgroup_vote"] - extensions = ["SPV_KHR_float_controls", "SPV_KHR_subgroup_vote"] +// CHECK: [Shader, Float16] +spv.module Logical GLSL450 requires #spv.vce { +} + +// ----- + +// CHECK: [SPV_KHR_float_controls, SPV_KHR_subgroup_vote] +spv.module Logical GLSL450 requires #spv.vce { } diff --git a/mlir/test/Dialect/SPIRV/Serialization/non-uniform-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/non-uniform-ops.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/non-uniform-ops.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/non-uniform-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @group_non_uniform_ballot spv.func @group_non_uniform_ballot(%predicate: i1) -> vector<4xi32> "None" { // CHECK: %{{.*}} = spv.GroupNonUniformBallot "Workgroup" %{{.*}}: vector<4xi32> diff --git a/mlir/test/Dialect/SPIRV/Serialization/phi.mlir b/mlir/test/Dialect/SPIRV/Serialization/phi.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/phi.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/phi.mlir @@ -2,7 +2,7 @@ // Test branch with one block argument -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { // CHECK: %[[CST:.*]] = spv.constant 0 %zero = spv.constant 0 : i32 @@ -17,15 +17,13 @@ spv.Return } spv.EntryPoint "GLCompute" @main -} attributes { - capabilities = ["Shader"] } // ----- // Test branch with multiple block arguments -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { // CHECK: %[[ZERO:.*]] = spv.constant 0 %zero = spv.constant 0 : i32 @@ -43,15 +41,13 @@ spv.Return } spv.EntryPoint "GLCompute" @main -} attributes { - capabilities = ["Shader"] } // ----- // Test using block arguments within branch -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { // CHECK: %[[CST0:.*]] = spv.constant 0 %zero = spv.constant 0 : i32 @@ -75,15 +71,13 @@ spv.Return } spv.EntryPoint "GLCompute" @main -} attributes { - capabilities = ["Shader"] } // ----- // Test block not following domination order -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { // CHECK: spv.Branch ^bb1 spv.Branch ^bb1 @@ -109,15 +103,13 @@ spv.Return } spv.EntryPoint "GLCompute" @main -} attributes { - capabilities = ["Shader"] } // ----- // Test multiple predecessors -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { %var = spv.Variable : !spv.ptr @@ -160,15 +152,13 @@ spv.Return } spv.EntryPoint "GLCompute" @main -} attributes { - capabilities = ["Shader"] } // ----- // Test nested loops with block arguments -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.globalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr, Input> spv.globalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr, Input> spv.func @fmul_kernel() "None" { @@ -245,4 +235,4 @@ spv.EntryPoint "GLCompute" @fmul_kernel, @__builtin_var_WorkgroupId__, @__builtin_var_NumWorkgroups__ spv.ExecutionMode @fmul_kernel "LocalSize", 32, 1, 1 -} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]} +} diff --git a/mlir/test/Dialect/SPIRV/Serialization/selection.mlir b/mlir/test/Dialect/SPIRV/Serialization/selection.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/selection.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/selection.mlir @@ -2,7 +2,7 @@ // Selection with both then and else branches -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @selection(%cond: i1) -> () "None" { // CHECK: spv.Branch ^bb1 // CHECK-NEXT: ^bb1: @@ -48,8 +48,6 @@ } spv.EntryPoint "GLCompute" @main spv.ExecutionMode @main "LocalSize", 1, 1, 1 -} attributes { - capabilities = ["Shader"] } // ----- @@ -57,7 +55,7 @@ // Selection with only then branch // Selection in function entry block -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK: spv.func @selection(%[[ARG:.*]]: i1 spv.func @selection(%cond: i1) -> (i32) "None" { // CHECK: spv.Branch ^bb1 @@ -87,7 +85,5 @@ } spv.EntryPoint "GLCompute" @main spv.ExecutionMode @main "LocalSize", 1, 1, 1 -} attributes { - capabilities = ["Shader"] } diff --git a/mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir b/mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK: spv.specConstant @sc_true = true spv.specConstant @sc_true = true // CHECK: spv.specConstant @sc_false spec_id(1) = false diff --git a/mlir/test/Dialect/SPIRV/Serialization/struct.mlir b/mlir/test/Dialect/SPIRV/Serialization/struct.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/struct.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/struct.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK: !spv.ptr [0]>, Input> spv.globalVariable @var0 bind(0, 1) : !spv.ptr [0]>, Input> diff --git a/mlir/test/Dialect/SPIRV/Serialization/terminator.mlir b/mlir/test/Dialect/SPIRV/Serialization/terminator.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/terminator.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/terminator.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @ret spv.func @ret() -> () "None" { // CHECK: spv.Return diff --git a/mlir/test/Dialect/SPIRV/Serialization/undef.mlir b/mlir/test/Dialect/SPIRV/Serialization/undef.mlir --- a/mlir/test/Dialect/SPIRV/Serialization/undef.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/undef.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { // CHECK: {{%.*}} = spv.undef : f32 // CHECK-NEXT: {{%.*}} = spv.undef : f32 @@ -23,7 +23,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 requires #spv.vce { // CHECK: spv.func {{@.*}} spv.func @ignore_unused_undef() -> () "None" { // CHECK-NEXT: spv.Return diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir @@ -1,7 +1,7 @@ // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s // CHECK-LABEL: spv.module -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // CHECK-DAG: spv.globalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") spv.globalVariable @__builtin_var_WorkgroupSize__ built_in("WorkgroupSize") : !spv.ptr, Input> // CHECK-DAG: spv.globalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") @@ -122,4 +122,4 @@ } // CHECK: spv.EntryPoint "GLCompute" [[FN]], [[WORKGROUPID]], [[LOCALINVOCATIONID]], [[NUMWORKGROUPS]], [[WORKGROUPSIZE]] // CHECK-NEXT: spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1 -} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]} +} diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-simple.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-simple.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/abi-simple.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/abi-simple.mlir @@ -1,7 +1,7 @@ // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s // CHECK-LABEL: spv.module -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // CHECK-DAG: spv.globalVariable [[VAR0:@.*]] bind(0, 0) : !spv.ptr, StorageBuffer> // CHECK-DAG: spv.globalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr [0]>, StorageBuffer> // CHECK: spv.func [[FN:@.*]]() @@ -24,4 +24,4 @@ } // CHECK: spv.EntryPoint "GLCompute" [[FN]] // CHECK: spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1 -} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]} +} diff --git a/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir b/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir @@ -1,6 +1,6 @@ // RUN: mlir-opt %s -split-input-file -pass-pipeline='spv.module(inline)' -mlir-disable-inline-simplify | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @callee() "None" { spv.Return } @@ -15,7 +15,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @callee() -> i32 "None" { %0 = spv.constant 42 : i32 spv.ReturnValue %0 : i32 @@ -32,7 +32,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.globalVariable @data bind(0, 0) : !spv.ptr [0]>, StorageBuffer> spv.func @callee() "None" { %0 = spv._address_of @data : !spv.ptr [0]>, StorageBuffer> @@ -67,7 +67,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @callee(%cond : i1) -> () "None" { spv.selection { spv.BranchConditional %cond, ^then, ^merge @@ -90,7 +90,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @callee(%cond : i1) -> () "None" { spv.selection { spv.BranchConditional %cond, ^then, ^merge @@ -119,7 +119,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @callee(%cond : i1) -> () "None" { spv.loop { spv.Branch ^header @@ -146,7 +146,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @callee(%cond : i1) -> () "None" { spv.loop { spv.Branch ^header @@ -183,7 +183,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.globalVariable @arg_0 bind(0, 0) : !spv.ptr, StorageBuffer> spv.globalVariable @arg_1 bind(0, 1) : !spv.ptr, StorageBuffer> @@ -222,7 +222,7 @@ } spv.EntryPoint "GLCompute" @inline_into_selection_region spv.ExecutionMode @inline_into_selection_region "LocalSize", 32, 1, 1 -} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]} +} // TODO: Add tests for inlining structured control flow into // structured control flow. diff --git a/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir b/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir @@ -1,6 +1,6 @@ // RUN: mlir-opt -decorate-spirv-composite-type-layout -split-input-file -verify-diagnostics %s -o - | FileCheck %s -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // CHECK: spv.globalVariable @var0 bind(0, 1) : !spv.ptr [4], f32 [12]>, Uniform> spv.globalVariable @var0 bind(0,1) : !spv.ptr, f32>, Uniform> @@ -31,7 +31,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // CHECK: spv.globalVariable @var0 : !spv.ptr [0], i1 [16]> [0], i1 [24]> [0], i1 [32]> [0], i1 [40]>, Uniform> spv.globalVariable @var0 : !spv.ptr, i1>, i1>, i1>, i1>, Uniform> @@ -59,7 +59,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // CHECK: spv.globalVariable @var0 : !spv.ptr [0], f32 [8]>, StorageBuffer> spv.globalVariable @var0 : !spv.ptr, f32>, StorageBuffer> @@ -72,7 +72,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // CHECK: spv.globalVariable @emptyStructAsMember : !spv.ptr [0]>, StorageBuffer> spv.globalVariable @emptyStructAsMember : !spv.ptr>, StorageBuffer> @@ -91,7 +91,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // CHECK: spv.globalVariable @var0 : !spv.ptr, PushConstant> spv.globalVariable @var0 : !spv.ptr, PushConstant> // CHECK: spv.globalVariable @var1 : !spv.ptr, PhysicalStorageBuffer> diff --git a/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir b/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir @@ -7,33 +7,33 @@ // Test deducing minimal version. // spv.IAdd is available from v1.0. -// CHECK: vce_triple = #spv.vce -spv.module "Logical" "GLSL450" { - spv.func @iadd(%val : i32) -> i32 "None" { - %0 = spv.IAdd %val, %val: i32 - spv.ReturnValue %0: i32 - } -} attributes { +// CHECK: requires #spv.vce +spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> +} { + spv.func @iadd(%val : i32) -> i32 "None" { + %0 = spv.IAdd %val, %val: i32 + spv.ReturnValue %0: i32 + } } // Test deducing minimal version. // spv.GroupNonUniformBallot is available since v1.3. -// CHECK: vce_triple = #spv.vce -spv.module "Logical" "GLSL450" { - spv.func @group_non_uniform_ballot(%predicate : i1) -> vector<4xi32> "None" { - %0 = spv.GroupNonUniformBallot "Workgroup" %predicate : vector<4xi32> - spv.ReturnValue %0: vector<4xi32> - } -} attributes { +// CHECK: requires #spv.vce +spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> +} { + spv.func @group_non_uniform_ballot(%predicate : i1) -> vector<4xi32> "None" { + %0 = spv.GroupNonUniformBallot "Workgroup" %predicate : vector<4xi32> + spv.ReturnValue %0: vector<4xi32> + } } //===----------------------------------------------------------------------===// @@ -42,33 +42,33 @@ // Test minimal capabilities. -// CHECK: vce_triple = #spv.vce -spv.module "Logical" "GLSL450" { - spv.func @iadd(%val : i32) -> i32 "None" { - %0 = spv.IAdd %val, %val: i32 - spv.ReturnValue %0: i32 - } -} attributes { +// CHECK: requires #spv.vce +spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> +} { + spv.func @iadd(%val : i32) -> i32 "None" { + %0 = spv.IAdd %val, %val: i32 + spv.ReturnValue %0: i32 + } } // Test deducing implied capability. // AtomicStorage implies Shader. -// CHECK: vce_triple = #spv.vce -spv.module "Logical" "GLSL450" { - spv.func @iadd(%val : i32) -> i32 "None" { - %0 = spv.IAdd %val, %val: i32 - spv.ReturnValue %0: i32 - } -} attributes { +// CHECK: requires #spv.vce +spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> +} { + spv.func @iadd(%val : i32) -> i32 "None" { + %0 = spv.IAdd %val, %val: i32 + spv.ReturnValue %0: i32 + } } // Test selecting the capability available in the target environment. @@ -81,30 +81,30 @@ // * GroupNonUniformArithmetic // * GroupNonUniformBallot -// CHECK: vce_triple = #spv.vce -spv.module "Logical" "GLSL450" { - spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" { - %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32 - spv.ReturnValue %0: i32 - } -} attributes { +// CHECK: requires #spv.vce +spv.module Logical GLSL450 attributes { 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: vce_triple = #spv.vce -spv.module "Logical" "GLSL450" { +} { spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" { %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32 spv.ReturnValue %0: i32 } -} attributes { +} + +// CHECK: requires #spv.vce +spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> +} { + spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" { + %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32 + spv.ReturnValue %0: i32 + } } //===----------------------------------------------------------------------===// @@ -114,33 +114,33 @@ // Test deducing minimal extensions. // spv.SubgroupBallotKHR requires the SPV_KHR_shader_ballot extension. -// CHECK: vce_triple = #spv.vce -spv.module "Logical" "GLSL450" { - spv.func @subgroup_ballot(%predicate : i1) -> vector<4xi32> "None" { - %0 = spv.SubgroupBallotKHR %predicate: vector<4xi32> - spv.ReturnValue %0: vector<4xi32> - } -} attributes { +// CHECK: requires #spv.vce +spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> +} { + spv.func @subgroup_ballot(%predicate : i1) -> vector<4xi32> "None" { + %0 = spv.SubgroupBallotKHR %predicate: vector<4xi32> + spv.ReturnValue %0: vector<4xi32> + } } // Test deducing implied extension. // Vulkan memory model requires SPV_KHR_vulkan_memory_model, which is enabled // implicitly by v1.5. -// CHECK: vce_triple = #spv.vce -spv.module "Logical" "Vulkan" { - spv.func @iadd(%val : i32) -> i32 "None" { - %0 = spv.IAdd %val, %val: i32 - spv.ReturnValue %0: i32 - } -} attributes { +// CHECK: requires #spv.vce +spv.module Logical Vulkan attributes { spv.target_env = #spv.target_env< #spv.vce, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> +} { + spv.func @iadd(%val : i32) -> i32 "None" { + %0 = spv.IAdd %val, %val: i32 + spv.ReturnValue %0: i32 + } } diff --git a/mlir/test/Dialect/SPIRV/availability.mlir b/mlir/test/Dialect/SPIRV/availability.mlir --- a/mlir/test/Dialect/SPIRV/availability.mlir +++ b/mlir/test/Dialect/SPIRV/availability.mlir @@ -36,7 +36,7 @@ // CHECK: spv.module max version: v1.5 // CHECK: spv.module extensions: [ ] // CHECK: spv.module capabilities: [ [Shader] ] - spv.module "Logical" "GLSL450" { } + spv.module Logical GLSL450 { } return } @@ -46,6 +46,6 @@ // CHECK: spv.module max version: v1.5 // CHECK: spv.module extensions: [ [SPV_EXT_physical_storage_buffer, SPV_KHR_physical_storage_buffer] [SPV_KHR_vulkan_memory_model] ] // CHECK: spv.module capabilities: [ [PhysicalStorageBufferAddresses] [VulkanMemoryModel] ] - spv.module "PhysicalStorageBuffer64" "Vulkan" { } + spv.module PhysicalStorageBuffer64 Vulkan { } return } diff --git a/mlir/test/Dialect/SPIRV/control-flow-ops.mlir b/mlir/test/Dialect/SPIRV/control-flow-ops.mlir --- a/mlir/test/Dialect/SPIRV/control-flow-ops.mlir +++ b/mlir/test/Dialect/SPIRV/control-flow-ops.mlir @@ -155,7 +155,7 @@ // spv.FunctionCall //===----------------------------------------------------------------------===// -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @fmain(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>, %arg2 : i32) -> i32 "None" { // CHECK: {{%.*}} = spv.FunctionCall @f_0({{%.*}}, {{%.*}}) : (vector<4xf32>, vector<4xf32>) -> vector<4xf32> %0 = spv.FunctionCall @f_0(%arg0, %arg1) : (vector<4xf32>, vector<4xf32>) -> vector<4xf32> @@ -200,7 +200,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @f_invalid_result_type(%arg0 : i32, %arg1 : i32) -> () "None" { // expected-error @+1 {{expected callee function to have 0 or 1 result, but provided 2}} %0:2 = spv.FunctionCall @f_invalid_result_type(%arg0, %arg1) : (i32, i32) -> (i32, i32) @@ -210,7 +210,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @f_result_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" { // expected-error @+1 {{has incorrect number of results has for callee: expected 0, but provided 1}} %1 = spv.FunctionCall @f_result_type_mismatch(%arg0, %arg0) : (i32, i32) -> (i32) @@ -220,7 +220,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" { // expected-error @+1 {{has incorrect number of operands for callee: expected 2, but provided 1}} spv.FunctionCall @f_type_mismatch(%arg0) : (i32) -> () @@ -230,7 +230,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" { %0 = spv.constant 2.0 : f32 // expected-error @+1 {{operand type mismatch: expected operand type 'i32', but provided 'f32' for operand number 1}} @@ -241,7 +241,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> i32 "None" { %cst = spv.constant 0: i32 // expected-error @+1 {{result type mismatch: expected 'i32', but provided 'f32'}} @@ -252,7 +252,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @f_foo(%arg0 : i32, %arg1 : i32) -> i32 "None" { // expected-error @+1 {{op callee function 'f_undefined' not found in nearest symbol table}} %0 = spv.FunctionCall @f_undefined(%arg0, %arg0) : (i32, i32) -> i32 @@ -518,7 +518,7 @@ // ----- // Return mismatches function signature -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @work() -> (i32) "None" { // expected-error @+1 {{cannot be used in functions returning value}} spv.Return @@ -527,7 +527,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @in_nested_region(%cond: i1) -> (i32) "None" { spv.selection { spv.BranchConditional %cond, ^then, ^merge @@ -605,7 +605,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @value_count_mismatch() -> () "None" { %0 = spv.constant 42 : i32 // expected-error @+1 {{op returns 1 value but enclosing function requires 0 results}} @@ -615,7 +615,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @value_type_mismatch() -> (f32) "None" { %0 = spv.constant 42 : i32 // expected-error @+1 {{return value's type ('i32') mismatch with function's result type ('f32')}} @@ -625,7 +625,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @in_nested_region(%cond: i1) -> () "None" { spv.selection { spv.BranchConditional %cond, ^then, ^merge diff --git a/mlir/test/Dialect/SPIRV/ops.mlir b/mlir/test/Dialect/SPIRV/ops.mlir --- a/mlir/test/Dialect/SPIRV/ops.mlir +++ b/mlir/test/Dialect/SPIRV/ops.mlir @@ -416,7 +416,7 @@ // spv.ExecutionMode //===----------------------------------------------------------------------===// -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -425,7 +425,7 @@ spv.ExecutionMode @do_nothing "ContractionOff" } -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -436,7 +436,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -639,7 +639,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.globalVariable @var0 : !spv.ptr // CHECK_LABEL: @simple_load spv.func @simple_load() -> () "None" { @@ -1057,7 +1057,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.globalVariable @var0 : !spv.ptr spv.func @simple_store(%arg0 : f32) -> () "None" { %0 = spv._address_of @var0 : !spv.ptr @@ -1130,7 +1130,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.globalVariable @global : !spv.ptr spv.func @variable_init_global_variable() -> () "None" { %0 = spv._address_of @global : !spv.ptr @@ -1138,14 +1138,11 @@ %1 = spv.Variable init(%0) : !spv.ptr, Function> spv.Return } -} attributes { - capability = ["VariablePointers"], - extension = ["SPV_KHR_variable_pointers"] } // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.specConstant @sc = 42 : i32 // CHECK-LABEL: @variable_init_spec_constant spv.func @variable_init_spec_constant() -> () "None" { diff --git a/mlir/test/Dialect/SPIRV/structure-ops.mlir b/mlir/test/Dialect/SPIRV/structure-ops.mlir --- a/mlir/test/Dialect/SPIRV/structure-ops.mlir +++ b/mlir/test/Dialect/SPIRV/structure-ops.mlir @@ -4,7 +4,7 @@ // spv._address_of //===----------------------------------------------------------------------===// -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.globalVariable @var1 : !spv.ptr>, Input> spv.func @access_chain() -> () "None" { %0 = spv.constant 1: i32 @@ -28,7 +28,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.globalVariable @var1 : !spv.ptr>, Input> spv.func @foo() -> () "None" { // expected-error @+1 {{expected spv.globalVariable symbol}} @@ -38,7 +38,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.globalVariable @var1 : !spv.ptr>, Input> spv.func @foo() -> () "None" { // expected-error @+1 {{result type mismatch with the referenced global variable's type}} @@ -135,7 +135,7 @@ // spv.EntryPoint //===----------------------------------------------------------------------===// -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -143,7 +143,7 @@ spv.EntryPoint "GLCompute" @do_nothing } -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.globalVariable @var2 : !spv.ptr spv.globalVariable @var3 : !spv.ptr spv.func @do_something(%arg0 : !spv.ptr, %arg1 : !spv.ptr) -> () "None" { @@ -157,7 +157,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -167,7 +167,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -182,7 +182,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { // expected-error @+1 {{op must appear in a module-like op's block}} spv.EntryPoint "GLCompute" @do_something @@ -191,7 +191,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -202,12 +202,12 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } spv.EntryPoint "GLCompute" @do_nothing - // expected-error @+1 {{custom op 'spv.EntryPoint' invalid execution_model attribute specification: "ContractionOff"}} + // expected-error @+1 {{'spv.EntryPoint' invalid execution_model attribute specification: "ContractionOff"}} spv.EntryPoint "ContractionOff" @do_nothing } @@ -250,7 +250,7 @@ // ----- // Nested function -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @outer_func() -> () "None" { // expected-error @+1 {{must appear in a module-like op's block}} spv.func @inner_func() -> () "None" { @@ -266,13 +266,13 @@ // spv.globalVariable //===----------------------------------------------------------------------===// -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // CHECK: spv.globalVariable @var0 : !spv.ptr spv.globalVariable @var0 : !spv.ptr } // TODO: Fix test case after initialization with normal constant is addressed -// spv.module "Logical" "GLSL450" { +// spv.module Logical GLSL450 { // %0 = spv.constant 4.0 : f32 // // CHECK1: spv.Variable init(%0) : !spv.ptr // spv.globalVariable @var1 init(%0) : !spv.ptr @@ -280,7 +280,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.specConstant @sc = 4.0 : f32 // CHECK: spv.globalVariable @var initializer(@sc) : !spv.ptr spv.globalVariable @var initializer(@sc) : !spv.ptr @@ -295,13 +295,13 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // CHECK: spv.globalVariable @var0 bind(1, 2) : !spv.ptr spv.globalVariable @var0 bind(1, 2) : !spv.ptr } // TODO: Fix test case after initialization with constant is addressed -// spv.module "Logical" "GLSL450" { +// spv.module Logical GLSL450 { // %0 = spv.constant 4.0 : f32 // // CHECK1: spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr // spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr @@ -309,7 +309,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // CHECK: spv.globalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr, Input> spv.globalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr, Input> // CHECK: spv.globalVariable @var2 built_in("GlobalInvocationID") : !spv.ptr, Input> @@ -326,28 +326,28 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // expected-error @+1 {{expected spv.ptr type}} spv.globalVariable @var0 : f32 } // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // expected-error @+1 {{op initializer must be result of a spv.specConstant or spv.globalVariable op}} spv.globalVariable @var0 initializer(@var1) : !spv.ptr } // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // expected-error @+1 {{storage class cannot be 'Generic'}} spv.globalVariable @var0 : !spv.ptr } // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @foo() "None" { // expected-error @+1 {{op must appear in a module-like op's block}} spv.globalVariable @var0 : !spv.ptr @@ -362,25 +362,33 @@ //===----------------------------------------------------------------------===// // Module without capability and extension -// CHECK: spv.module "Logical" "GLSL450" -spv.module "Logical" "GLSL450" { } +// CHECK: spv.module Logical GLSL450 +spv.module Logical GLSL450 { } -// Module with capability and extension -// CHECK: attributes {capability = ["Shader"], extension = ["SPV_KHR_16bit_storage"]} -spv.module "Logical" "GLSL450" { } attributes { - capability = ["Shader"], - extension = ["SPV_KHR_16bit_storage"] -} + +// Module with (version, capabilities, extensions) triple +// CHECK: spv.module Logical GLSL450 requires #spv.vce +spv.module Logical GLSL450 requires #spv.vce { } + +// Module with additional attributes +// CHECK: spv.module Logical GLSL450 attributes {foo = "bar"} +spv.module Logical GLSL450 attributes {foo = "bar"} { } + +// Module with VCE triple and additional attributes +// CHECK: spv.module Logical GLSL450 requires #spv.vce attributes {foo = "bar"} +spv.module Logical GLSL450 + requires #spv.vce + attributes {foo = "bar"} { } // Module with explicit spv._module_end // CHECK: spv.module -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv._module_end } // Module with function // CHECK: spv.module -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -389,32 +397,32 @@ // ----- // Missing addressing model -// expected-error@+1 {{custom op 'spv.module' expected addressing_model attribute specified as string}} +// expected-error@+1 {{'spv.module' expected valid keyword}} spv.module { } // ----- // Wrong addressing model -// expected-error@+1 {{custom op 'spv.module' invalid addressing_model attribute specification: "Physical"}} -spv.module "Physical" { } +// expected-error@+1 {{'spv.module' invalid addressing_model attribute specification: Physical}} +spv.module Physical { } // ----- // Missing memory model -// expected-error@+1 {{custom op 'spv.module' expected memory_model attribute specified as string}} -spv.module "Logical" { } +// expected-error@+1 {{'spv.module' expected valid keyword}} +spv.module Logical { } // ----- // Wrong memory model -// expected-error@+1 {{custom op 'spv.module' invalid memory_model attribute specification: "Bla"}} -spv.module "Logical" "Bla" { } +// expected-error@+1 {{'spv.module' invalid memory_model attribute specification: Bla}} +spv.module Logical Bla { } // ----- // Module with multiple blocks // expected-error @+1 {{expects region #0 to have 0 or 1 blocks}} -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { ^first: spv.Return ^second: @@ -433,7 +441,7 @@ // ----- // Use non SPIR-V op inside module -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // expected-error @+1 {{'spv.module' can only contain spv.* ops}} "dialect.op"() : () -> () } @@ -441,7 +449,7 @@ // ----- // Use non SPIR-V op inside function -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { // expected-error @+1 {{functions in 'spv.module' can only contain spv.* ops}} "dialect.op"() : () -> () @@ -451,29 +459,13 @@ // ----- // Use external function -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // expected-error @+1 {{'spv.module' cannot contain external functions}} spv.func @extern() -> () "None" } // ----- -// expected-error @+1 {{uses unknown capability: MyAwesomeCapability}} -spv.module "Logical" "GLSL450" { -} attributes { - capabilities = ["MyAwesomeCapability"] -} - -// ----- - -// expected-error @+1 {{uses unknown extension: MyAwesomeExtension}} -spv.module "Logical" "GLSL450" { -} attributes { - extensions = ["MyAwesomeExtension"] -} - -// ----- - //===----------------------------------------------------------------------===// // spv._module_end //===----------------------------------------------------------------------===// @@ -489,7 +481,7 @@ // spv._reference_of //===----------------------------------------------------------------------===// -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.specConstant @sc1 = false spv.specConstant @sc2 = 42 : i64 spv.specConstant @sc3 = 1.5 : f32 @@ -532,7 +524,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.func @foo() -> () "None" { // expected-error @+1 {{expected spv.specConstant symbol}} %0 = spv._reference_of @sc : i32 @@ -542,7 +534,7 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { spv.specConstant @sc = 42 : i32 spv.func @foo() -> () "None" { // expected-error @+1 {{result type mismatch with the referenced specialization constant's type}} @@ -557,7 +549,7 @@ // spv.specConstant //===----------------------------------------------------------------------===// -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // CHECK: spv.specConstant @sc1 = false spv.specConstant @sc1 = false // CHECK: spv.specConstant @sc2 spec_id(5) = 42 : i64 @@ -568,21 +560,21 @@ // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // expected-error @+1 {{SpecId cannot be negative}} spv.specConstant @sc2 spec_id(-5) = 42 : i64 } // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // expected-error @+1 {{default value bitwidth disallowed}} spv.specConstant @sc = 15 : i4 } // ----- -spv.module "Logical" "GLSL450" { +spv.module Logical GLSL450 { // expected-error @+1 {{default value can only be a bool, integer, or float scalar}} spv.specConstant @sc = dense<[2, 3]> : vector<2xi32> } diff --git a/mlir/test/Dialect/SPIRV/target-env.mlir b/mlir/test/Dialect/SPIRV/target-env.mlir --- a/mlir/test/Dialect/SPIRV/target-env.mlir +++ b/mlir/test/Dialect/SPIRV/target-env.mlir @@ -148,7 +148,7 @@ func @module_suitable_extension1() attributes { 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" + // CHECK: spv.module PhysicalStorageBuffer64 Vulkan "test.convert_to_module_op"() : () ->() return } @@ -157,7 +157,7 @@ func @module_suitable_extension2() attributes { 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" + // CHECK: spv.module PhysicalStorageBuffer64 Vulkan "test.convert_to_module_op"() : () -> () return } @@ -185,7 +185,7 @@ // Version 1.5 implies SPV_KHR_vulkan_memory_model and SPV_KHR_physical_storage_buffer. 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" + // CHECK: spv.module PhysicalStorageBuffer64 Vulkan "test.convert_to_module_op"() : () -> () return } diff --git a/mlir/test/mlir-vulkan-runner/addf.mlir b/mlir/test/mlir-vulkan-runner/addf.mlir --- a/mlir/test/mlir-vulkan-runner/addf.mlir +++ b/mlir/test/mlir-vulkan-runner/addf.mlir @@ -1,7 +1,13 @@ // RUN: mlir-vulkan-runner %s --shared-libs=%vulkan_wrapper_library_dir/libvulkan-runtime-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s // CHECK: [3.3, 3.3, 3.3, 3.3, 3.3, 3.3, 3.3, 3.3] -module attributes {gpu.container_module} { +module attributes { + gpu.container_module, + spv.target_env = #spv.target_env< + #spv.vce, + {max_compute_workgroup_invocations = 128 : i32, + max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> +} { gpu.module @kernels { gpu.func @kernel_add(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>) attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32>}} { diff --git a/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp b/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp --- a/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp +++ b/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp @@ -38,6 +38,7 @@ passManager.addPass(createConvertGPUToSPIRVPass()); OpPassManager &modulePM = passManager.nest(); modulePM.addPass(spirv::createLowerABIAttributesPass()); + modulePM.addPass(spirv::createUpdateVersionCapabilityExtensionPass()); passManager.addPass(createConvertGpuLaunchFuncToVulkanLaunchFuncPass()); passManager.addPass(createLowerToLLVMPass()); passManager.addPass(createConvertVulkanLaunchFuncToVulkanCallsPass()); diff --git a/mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp b/mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp --- a/mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp +++ b/mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp @@ -63,7 +63,9 @@ //===--------------------------------------------------------------------===// /// Adds the SPIR-V module header to `binary`. - void addHeader() { spirv::appendModuleHeader(binary, /*idBound=*/0); } + void addHeader() { + spirv::appendModuleHeader(binary, spirv::Version::V_1_0, /*idBound=*/0); + } /// Adds the SPIR-V instruction into `binary`. void addInstruction(spirv::Opcode op, ArrayRef operands) { diff --git a/mlir/unittests/Dialect/SPIRV/SerializationTest.cpp b/mlir/unittests/Dialect/SPIRV/SerializationTest.cpp --- a/mlir/unittests/Dialect/SPIRV/SerializationTest.cpp +++ b/mlir/unittests/Dialect/SPIRV/SerializationTest.cpp @@ -12,6 +12,7 @@ //===----------------------------------------------------------------------===// #include "mlir/Dialect/SPIRV/Serialization.h" +#include "mlir/Dialect/SPIRV/SPIRVAttributes.h" #include "mlir/Dialect/SPIRV/SPIRVBinaryUtils.h" #include "mlir/Dialect/SPIRV/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/SPIRVOps.h" @@ -46,6 +47,10 @@ state.addAttribute("memory_model", builder.getI32IntegerAttr( static_cast(spirv::MemoryModel::GLSL450))); + state.addAttribute("vce_triple", + spirv::VerCapExtAttr::get( + spirv::Version::V_1_0, ArrayRef(), + ArrayRef(), &context)); spirv::ModuleOp::build(&builder, state); module = cast(Operation::create(state)); }