[Mlir-commits] [mlir] 3148f10 - [mlir][spirv] Use spv.vce in spv.module and wire up (de)serialization
Lei Zhang
llvmlistbot at llvm.org
Thu Mar 12 16:39:56 PDT 2020
Author: Lei Zhang
Date: 2020-03-12T19:37:45-04:00
New Revision: 3148f10b1791e3227e194b1c974bb3e83a1c4d0e
URL: https://github.com/llvm/llvm-project/commit/3148f10b1791e3227e194b1c974bb3e83a1c4d0e
DIFF: https://github.com/llvm/llvm-project/commit/3148f10b1791e3227e194b1c974bb3e83a1c4d0e.diff
LOG: [mlir][spirv] Use spv.vce in spv.module and wire up (de)serialization
This commits changes the definition of spv.module to use the #spv.vce
attribute for specifying (version, capabilities, extensions) triple
so that we can have better API and custom assembly form. Since now
we have proper modelling of the triple, (de)serialization is wired up
to use them.
With the new UpdateVCEPass, we don't need to manually specify the
required extensions and capabilities anymore when creating a spv.module.
One just need to call UpdateVCEPass before serialization to get the
needed version/extensions/capabilities.
Differential Revision: https://reviews.llvm.org/D75872
Added:
Modified:
mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td
mlir/include/mlir/Dialect/SPIRV/SPIRVBinaryUtils.h
mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h
mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td
mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp
mlir/lib/Dialect/SPIRV/SPIRVOps.cpp
mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp
mlir/lib/Dialect/SPIRV/Serialization/SPIRVBinaryUtils.cpp
mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp
mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp
mlir/test/Conversion/GPUToSPIRV/builtins.mlir
mlir/test/Conversion/GPUToSPIRV/load-store.mlir
mlir/test/Conversion/GPUToSPIRV/simple.mlir
mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
mlir/test/Dialect/SPIRV/Serialization/arithmetic-ops.mlir
mlir/test/Dialect/SPIRV/Serialization/array.mlir
mlir/test/Dialect/SPIRV/Serialization/atomic-ops.mlir
mlir/test/Dialect/SPIRV/Serialization/barrier.mlir
mlir/test/Dialect/SPIRV/Serialization/bit-ops.mlir
mlir/test/Dialect/SPIRV/Serialization/cast-ops.mlir
mlir/test/Dialect/SPIRV/Serialization/composite-op.mlir
mlir/test/Dialect/SPIRV/Serialization/constant.mlir
mlir/test/Dialect/SPIRV/Serialization/entry-point.mlir
mlir/test/Dialect/SPIRV/Serialization/execution-mode.mlir
mlir/test/Dialect/SPIRV/Serialization/function-call.mlir
mlir/test/Dialect/SPIRV/Serialization/global-variable.mlir
mlir/test/Dialect/SPIRV/Serialization/glsl-ops.mlir
mlir/test/Dialect/SPIRV/Serialization/group-ops.mlir
mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir
mlir/test/Dialect/SPIRV/Serialization/loop.mlir
mlir/test/Dialect/SPIRV/Serialization/memory-ops.mlir
mlir/test/Dialect/SPIRV/Serialization/module.mlir
mlir/test/Dialect/SPIRV/Serialization/non-uniform-ops.mlir
mlir/test/Dialect/SPIRV/Serialization/phi.mlir
mlir/test/Dialect/SPIRV/Serialization/selection.mlir
mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir
mlir/test/Dialect/SPIRV/Serialization/struct.mlir
mlir/test/Dialect/SPIRV/Serialization/terminator.mlir
mlir/test/Dialect/SPIRV/Serialization/undef.mlir
mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir
mlir/test/Dialect/SPIRV/Transforms/abi-simple.mlir
mlir/test/Dialect/SPIRV/Transforms/inlining.mlir
mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir
mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir
mlir/test/Dialect/SPIRV/availability.mlir
mlir/test/Dialect/SPIRV/control-flow-ops.mlir
mlir/test/Dialect/SPIRV/ops.mlir
mlir/test/Dialect/SPIRV/structure-ops.mlir
mlir/test/Dialect/SPIRV/target-env.mlir
mlir/test/mlir-vulkan-runner/addf.mlir
mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp
mlir/unittests/Dialect/SPIRV/SerializationTest.cpp
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td
index 1af6ddef4ea0..8ef1e363eebc 100644
--- a/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td
+++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td
@@ -2941,6 +2941,18 @@ def SPV_SamplerUseAttr:
"ImageSamplerUseInfo", "valid SPIR-V Sampler Use specification",
[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
index 0c19c6783ed1..913ba4193b7c 100644
--- a/mlir/include/mlir/Dialect/SPIRV/SPIRVBinaryUtils.h
+++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVBinaryUtils.h
@@ -34,8 +34,10 @@ constexpr uint32_t kGeneratorNumber = 22;
#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<uint32_t> &header, uint32_t idBound);
+/// Appends a SPRI-V module header to `header` with the given `version` and
+/// `idBound`.
+void appendModuleHeader(SmallVectorImpl<uint32_t> &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
index d4c7a1bd824d..9bb0d52095d4 100644
--- a/mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h
+++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h
@@ -23,6 +23,7 @@ namespace mlir {
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
index 26f8510a718d..ed8f0b1d3883 100644
--- a/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td
+++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td
@@ -382,25 +382,25 @@ def SPV_ModuleOp : SPV_Op<"module",
### 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<v1.0, [Shader], [SPV_KHR_vulkan_memory_model]>
+ attributes { some_additional_attr = ... } {
+ spv.func @do_nothing() -> () {
spv.Return
}
- } attributes {
- capability = ["Shader"],
- extension = ["SPV_KHR_16bit_storage"]
}
```
}];
@@ -408,26 +408,19 @@ def SPV_ModuleOp : SPV_Op<"module",
let arguments = (ins
SPV_AddressingModelAttr:$addressing_model,
SPV_MemoryModelAttr:$memory_model,
- OptionalAttr<StrArrayAttr>:$capabilities,
- OptionalAttr<StrArrayAttr>:$extensions,
- OptionalAttr<StrArrayAttr>:$extended_instruction_sets
+ OptionalAttr<SPV_VerCapExtAttr>:$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<spirv::Capability> capabilities = {},
- /*optional*/ ArrayRef<spirv::Extension> 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 @@ def SPV_ModuleOp : SPV_Op<"module",
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
index 110f8c5c2f59..3c07097db542 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp
@@ -376,13 +376,10 @@ PatternMatchResult GPUFuncOpConversion::matchAndRewrite(
PatternMatchResult GPUModuleConversion::matchAndRewrite(
gpu::GPUModuleOp moduleOp, ArrayRef<Value> operands,
ConversionPatternRewriter &rewriter) const {
- // TODO : Generalize this to account for
diff erent extensions,
- // capabilities, extended_instruction_sets, other addressing models
- // and memory models.
auto spvModule = rewriter.create<spirv::ModuleOp>(
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
index 25da3b80abfc..377242482b2a 100644
--- 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 @@ getStrArrayAttrForEnumList(Builder &builder, ArrayRef<Ty> enumValues,
return builder.getStrArrayAttr(enumValStrs);
}
+/// Parses the next string attribute in `parser` as an enumerant of the given
+/// `EnumClass`.
template <typename EnumClass>
static ParseResult
-parseEnumAttribute(EnumClass &value, OpAsmParser &parser,
- StringRef attrName = spirv::attributeName<EnumClass>()) {
+parseEnumStrAttr(EnumClass &value, OpAsmParser &parser,
+ StringRef attrName = spirv::attributeName<EnumClass>()) {
Attribute attrVal;
SmallVector<NamedAttribute, 1> attr;
auto loc = parser.getCurrentLocation();
@@ -122,11 +125,49 @@ parseEnumAttribute(EnumClass &value, OpAsmParser &parser,
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 <typename EnumClass>
static ParseResult
-parseEnumAttribute(EnumClass &value, OpAsmParser &parser, OperationState &state,
- StringRef attrName = spirv::attributeName<EnumClass>()) {
- if (parseEnumAttribute(value, parser)) {
+parseEnumStrAttr(EnumClass &value, OpAsmParser &parser, OperationState &state,
+ StringRef attrName = spirv::attributeName<EnumClass>()) {
+ if (parseEnumStrAttr(value, parser)) {
+ return failure();
+ }
+ state.addAttribute(attrName, parser.getBuilder().getI32IntegerAttr(
+ llvm::bit_cast<int32_t>(value)));
+ return success();
+}
+
+/// Parses the next keyword in `parser` as an enumerant of the given
+/// `EnumClass`.
+template <typename EnumClass>
+static ParseResult
+parseEnumKeywordAttr(EnumClass &value, OpAsmParser &parser,
+ StringRef attrName = spirv::attributeName<EnumClass>()) {
+ StringRef keyword;
+ SmallVector<NamedAttribute, 1> attr;
+ auto loc = parser.getCurrentLocation();
+ if (parser.parseKeyword(&keyword))
+ return failure();
+ if (Optional<EnumClass> attr = spirv::symbolizeEnum<EnumClass>()(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 <typename EnumClass>
+static ParseResult
+parseEnumKeywordAttr(EnumClass &value, OpAsmParser &parser,
+ OperationState &state,
+ StringRef attrName = spirv::attributeName<EnumClass>()) {
+ if (parseEnumKeywordAttr(value, parser)) {
return failure();
}
state.addAttribute(attrName, parser.getBuilder().getI32IntegerAttr(
@@ -143,7 +184,7 @@ static ParseResult parseMemoryAccessAttributes(OpAsmParser &parser,
}
spirv::MemoryAccess memoryAccessAttr;
- if (parseEnumAttribute(memoryAccessAttr, parser, state)) {
+ if (parseEnumStrAttr(memoryAccessAttr, parser, state)) {
return failure();
}
@@ -463,8 +504,8 @@ static ParseResult parseAtomicUpdateOp(OpAsmParser &parser,
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 @@ static ParseResult parseGroupNonUniformArithmeticOp(OpAsmParser &parser,
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 @@ static ParseResult parseAtomicCompareExchangeWeakOp(OpAsmParser &parser,
spirv::MemorySemantics equalSemantics, unequalSemantics;
SmallVector<OpAsmParser::OperandType, 3> 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 @@ static ParseResult parseEntryPointOp(OpAsmParser &parser,
SmallVector<Attribute, 4> 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 @@ static ParseResult parseExecutionModeOp(OpAsmParser &parser,
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 @@ static ParseResult parseFuncOp(OpAsmParser &parser, OperationState &state) {
// 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 @@ static ParseResult parseLoadOp(OpAsmParser &parser, OperationState &state) {
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 @@ void spirv::ModuleOp::build(Builder *builder, OperationState &state) {
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<spirv::Capability> capabilities,
- ArrayRef<spirv::Extension> extensions,
- ArrayAttr extended_instruction_sets) {
+ spirv::MemoryModel memory_model) {
state.addAttribute(
"addressing_model",
builder->getI32IntegerAttr(static_cast<int32_t>(addressing_model)));
state.addAttribute("memory_model", builder->getI32IntegerAttr(
static_cast<int32_t>(memory_model)));
- if (!capabilities.empty())
- state.addAttribute("capabilities",
- getStrArrayAttrForEnumList<spirv::Capability>(
- *builder, capabilities, spirv::stringifyCapability));
- if (!extensions.empty())
- state.addAttribute("extensions",
- getStrArrayAttrForEnumList<spirv::Extension>(
- *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 @@ static ParseResult parseModuleOp(OpAsmParser &parser, OperationState &state) {
// 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 ParseResult parseModuleOp(OpAsmParser &parser, OperationState &state) {
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<StringRef, 2> elidedAttrs;
+
+ printer << " " << spirv::stringifyAddressingModel(moduleOp.addressing_model())
+ << " " << spirv::stringifyMemoryModel(moduleOp.memory_model());
auto addressingModelAttrName = spirv::attributeName<spirv::AddressingModel>();
auto memoryModelAttrName = spirv::attributeName<spirv::MemoryModel>();
- 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<spirv::VerCapExtAttr> 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<std::pair<spirv::FuncOp, spirv::ExecutionModel>, 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 @@ static LogicalResult verify(spirv::ModuleOp moduleOp) {
}
}
- // Verify capabilities. ODS already guarantees that we have an array of
- // string attributes.
- if (auto caps = moduleOp.getAttrOfType<ArrayAttr>("capabilities")) {
- for (auto cap : caps.getValue()) {
- auto capStr = cap.cast<StringAttr>().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<ArrayAttr>("extensions")) {
- for (auto ext : exts.getValue()) {
- auto extStr = ext.cast<StringAttr>().getValue();
- if (!spirv::symbolizeExtension(extStr))
- return moduleOp.emitOpError("uses unknown extension: ") << extStr;
- }
- }
-
return success();
}
@@ -2479,7 +2480,7 @@ static ParseResult parseStoreOp(OpAsmParser &parser, OperationState &state) {
SmallVector<OpAsmParser::OperandType, 2> 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
index 8614df46cb97..fbba3595db8f 100644
--- 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 @@ class Deserializer {
/// in the deserializer.
LogicalResult processCapability(ArrayRef<uint32_t> 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<uint32_t> words);
@@ -117,8 +115,9 @@ class Deserializer {
/// bookkeeping in the deserializer.
LogicalResult processExtInstImport(ArrayRef<uint32_t> 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<uint32_t> operands);
@@ -397,11 +396,13 @@ class Deserializer {
OpBuilder opBuilder;
+ spirv::Version version;
+
/// The list of capabilities used by the module.
llvm::SmallSetVector<spirv::Capability, 4> capabilities;
/// The list of extensions used by the module.
- llvm::SmallSetVector<StringRef, 2> extensions;
+ llvm::SmallSetVector<spirv::Extension, 2> extensions;
// Result <id> to type mapping.
DenseMap<uint32_t, Type> typeMap;
@@ -507,9 +508,7 @@ LogicalResult Deserializer::deserialize() {
}
}
- // 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 @@ Optional<spirv::ModuleOp> Deserializer::collect() { return module; }
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<spirv::ModuleOp>(Operation::create(state));
}
@@ -539,6 +535,32 @@ LogicalResult Deserializer::processHeader() {
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 @@ LogicalResult Deserializer::processCapability(ArrayRef<uint32_t> operands) {
return success();
}
-void Deserializer::attachCapabilities() {
- if (capabilities.empty())
- return;
-
- SmallVector<StringRef, 2> 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<uint32_t> words) {
if (words.empty()) {
return emitError(
@@ -579,12 +587,14 @@ LogicalResult Deserializer::processExtension(ArrayRef<uint32_t> words) {
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 @@ LogicalResult Deserializer::processExtInstImport(ArrayRef<uint32_t> words) {
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<uint32_t> operands) {
diff --git a/mlir/lib/Dialect/SPIRV/Serialization/SPIRVBinaryUtils.cpp b/mlir/lib/Dialect/SPIRV/Serialization/SPIRVBinaryUtils.cpp
index d98c53279814..eabc410fb972 100644
--- 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<uint32_t> &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 @@ void spirv::appendModuleHeader(SmallVectorImpl<uint32_t> &header,
// | 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); // <id> 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
index befee16c2ce2..3d5837b11318 100644
--- 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 @@ void Serializer::collect(SmallVectorImpl<uint32_t> &binary) {
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 @@ uint32_t Serializer::getOrCreateFunctionID(StringRef fnName) {
}
void Serializer::processCapability() {
- auto caps = module.getAttrOfType<ArrayAttr>("capabilities");
- if (!caps)
- return;
-
- for (auto cap : caps.getValue()) {
- auto capStr = cap.cast<StringAttr>().getValue();
- auto capVal = spirv::symbolizeCapability(capStr);
+ for (auto cap : module.vce_triple()->getCapabilities())
encodeInstructionInto(capabilities, spirv::Opcode::OpCapability,
- {static_cast<uint32_t>(*capVal)});
- }
+ {static_cast<uint32_t>(cap)});
}
void Serializer::processExtension() {
- auto exts = module.getAttrOfType<ArrayAttr>("extensions");
- if (!exts)
- return;
-
- SmallVector<uint32_t, 16> extName;
- for (auto ext : exts.getValue()) {
- auto extStr = ext.cast<StringAttr>().getValue();
+ llvm::SmallVector<uint32_t, 16> 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 Serializer::emitDecoration(uint32_t target,
LogicalResult spirv::serialize(spirv::ModuleOp module,
SmallVectorImpl<uint32_t> &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
index 26597dc46340..6647431b70fc 100644
--- a/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp
+++ b/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp
@@ -150,7 +150,7 @@ void UpdateVCEPass::runOnOperation() {
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<OpPassBase<spirv::ModuleOp>>
diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
index a3abd089d5af..e41002a71a0a 100644
--- a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
@@ -7,7 +7,7 @@ module attributes {gpu.container_module} {
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 @@ module attributes {gpu.container_module} {
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 @@ module attributes {gpu.container_module} {
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 @@ module attributes {gpu.container_module} {
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 @@ module attributes {gpu.container_module} {
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 @@ module attributes {gpu.container_module} {
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 @@ module attributes {gpu.container_module} {
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 @@ module attributes {gpu.container_module} {
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
index a99433aeb1f6..6588de870057 100644
--- a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir
@@ -15,7 +15,7 @@ module attributes {gpu.container_module} {
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<vector<3xi32>, Input>
// CHECK-DAG: spv.globalVariable [[LOCALINVOCATIONIDVAR:@.*]] built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
diff --git a/mlir/test/Conversion/GPUToSPIRV/simple.mlir b/mlir/test/Conversion/GPUToSPIRV/simple.mlir
index 8db63359ef0e..d9b32a6e571b 100644
--- 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<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
@@ -12,7 +12,6 @@ module attributes {gpu.container_module} {
// 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
index aa3daa04734e..c286b4c104db 100644
--- 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<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
spv.globalVariable @kernel_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer>
spv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
%0 = spv._address_of @kernel_arg_0 : !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer>
@@ -17,7 +17,7 @@ module attributes {gpu.container_module} {
}
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
index 47ab01e65392..55c67dafe6bb 100644
--- 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<v1.0, [Shader], []> {
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
index a2c70a84158a..aa7cc405b5ee 100644
--- 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<v1.0, [Shader], []> {
spv.func @array_stride(%arg0 : !spv.ptr<!spv.array<4x!spv.array<4xf32 [4]> [128]>, StorageBuffer>, %arg1 : i32, %arg2 : i32) "None" {
// CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.array<4 x !spv.array<4 x f32 [4]> [128]>, StorageBuffer>
%2 = spv.AccessChain %arg0[%arg1, %arg2] : !spv.ptr<!spv.array<4x!spv.array<4xf32 [4]> [128]>, StorageBuffer>
@@ -10,7 +10,7 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK: spv.globalVariable {{@.*}} : !spv.ptr<!spv.rtarray<f32>, StorageBuffer>
spv.globalVariable @var0 : !spv.ptr<!spv.rtarray<f32>, StorageBuffer>
// CHECK: spv.globalVariable {{@.*}} : !spv.ptr<!spv.rtarray<vector<4xf16>>, Input>
diff --git a/mlir/test/Dialect/SPIRV/Serialization/atomic-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/atomic-ops.mlir
index 3247d89bf3b7..6bf32af37155 100644
--- 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<v1.0, [Shader], []> {
// CHECK-LABEL: @atomic_compare_exchange_weak
spv.func @atomic_compare_exchange_weak(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 "None" {
// CHECK: spv.AtomicCompareExchangeWeak "Workgroup" "Release" "Acquire" %{{.*}}, %{{.*}}, %{{.*}} : !spv.ptr<i32, Workgroup>
diff --git a/mlir/test/Dialect/SPIRV/Serialization/barrier.mlir b/mlir/test/Dialect/SPIRV/Serialization/barrier.mlir
index 4a12b7c0e9c4..4c5735d86876 100644
--- 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<v1.0, [Shader], []> {
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
index bc959bf9cc2c..23bf788b20ce 100644
--- 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<v1.0, [Shader], []> {
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
index 7de4b5c6af2b..76bac23e6f8f 100644
--- 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<v1.0, [Shader], []> {
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" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
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
index 2f6227dca970..f6b7a4a54b67 100644
--- 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<v1.0, [Shader], []> {
spv.func @composite_insert(%arg0 : !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>>, %arg1: !spv.array<4xf32>) -> !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>> "None" {
// CHECK: spv.CompositeInsert {{%.*}}, {{%.*}}[1 : i32, 0 : i32] : !spv.array<4 x f32> into !spv.struct<f32, !spv.struct<!spv.array<4 x f32>, f32>>
%0 = spv.CompositeInsert %arg1, %arg0[1 : i32, 0 : i32] : !spv.array<4xf32> into !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>>
diff --git a/mlir/test/Dialect/SPIRV/Serialization/constant.mlir b/mlir/test/Dialect/SPIRV/Serialization/constant.mlir
index 86238728e98a..180bd2b644be 100644
--- 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<v1.0, [Shader], []> {
// 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
index 4a01cf44b53a..698c873e990a 100644
--- 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<v1.0, [Shader], []> {
spv.func @noop() -> () "None" {
spv.Return
}
@@ -12,7 +12,7 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK: spv.globalVariable @var2 : !spv.ptr<f32, Input>
// CHECK-NEXT: spv.globalVariable @var3 : !spv.ptr<f32, Output>
// CHECK-NEXT: spv.func @noop({{%.*}}: !spv.ptr<f32, Input>, {{%.*}}: !spv.ptr<f32, Output>) "None"
diff --git a/mlir/test/Dialect/SPIRV/Serialization/execution-mode.mlir b/mlir/test/Dialect/SPIRV/Serialization/execution-mode.mlir
index 77d17b61d219..1734cbf85472 100644
--- 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<v1.0, [Shader], []> {
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
index 4c2c4335c45d..e30b27e7c516 100644
--- 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<v1.0, [Shader], []> {
spv.globalVariable @var1 : !spv.ptr<!spv.array<4xf32>, 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
index c174b4318949..faa371ea9016 100644
--- 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<vector<3xi32>, Input>
// CHECK-NEXT: spv.globalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.globalVariable @var0 bind(1, 0) : !spv.ptr<f32, Input>
spv.globalVariable @var1 bind(0, 1) : !spv.ptr<f32, Output>
spv.globalVariable @var2 {built_in = "GlobalInvocationId"} : !spv.ptr<vector<3xi32>, Input>
@@ -14,7 +14,7 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK: spv.globalVariable @var1 : !spv.ptr<f32, Input>
// CHECK-NEXT: spv.globalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr<f32, Input>
spv.globalVariable @var1 : !spv.ptr<f32, Input>
@@ -23,7 +23,7 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
spv.func @foo() "None" {
// CHECK: %[[ADDR:.*]] = spv._address_of @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
diff --git a/mlir/test/Dialect/SPIRV/Serialization/glsl-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/glsl-ops.mlir
index e4da14b2a283..6cca9f556317 100644
--- 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<v1.0, [Shader], []> {
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
index 55647154a550..474e40b97acc 100644
--- 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<v1.0, [Shader], []> {
// 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
index 4269020a6a4d..77251e358741 100644
--- 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<v1.0, [Shader], []> {
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" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
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
index a22f433d8e6f..e280f21c38f1 100644
--- 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<v1.0, [Shader], []> {
// for (int i = 0; i < count; ++i) {}
spv.func @loop(%count : i32) -> () "None" {
%zero = spv.constant 0: i32
@@ -55,13 +55,11 @@ spv.module "Logical" "GLSL450" {
spv.Return
}
spv.EntryPoint "GLCompute" @main
-} attributes {
- capabilities = ["Shader"]
}
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.globalVariable @GV1 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
spv.globalVariable @GV2 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
spv.func @loop_kernel() "None" {
@@ -103,13 +101,13 @@ spv.module "Logical" "GLSL450" {
}
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<v1.0, [Shader], []> {
// for (int i = 0; i < count; ++i) {
// for (int j = 0; j < count; ++j) { }
// }
@@ -207,7 +205,5 @@ spv.module "Logical" "GLSL450" {
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
index d89f1fff2fc2..d082fa01d9ae 100644
--- 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<v1.0, [Shader], []> {
spv.func @load_store(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) "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" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @access_chain(%arg0 : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>, %arg1 : i32, %arg2 : i32) "None" {
// CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.array<4 x !spv.array<4 x f32>>, Function>
// CHECK-NEXT: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.array<4 x !spv.array<4 x f32>>, Function>
@@ -26,7 +26,7 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @load_store_zero_rank_float(%arg0: !spv.ptr<!spv.struct<!spv.array<1 x f32 [4]> [0]>, StorageBuffer>, %arg1: !spv.ptr<!spv.struct<!spv.array<1 x f32 [4]> [0]>, StorageBuffer>) "None" {
// CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<!spv.array<1 x f32 [4]> [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
index bf4c8111fe33..29973e9e4d77 100644
--- 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<v1.0, [Shader], []> {
// 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<v1.0, [Shader], []> {
spv.func @foo() -> () "None" {
spv.Return
}
@@ -14,17 +14,19 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
-} attributes {
- // CHECK: capabilities = ["Shader", "Float16"]
- capabilities = ["Shader", "Float16"]
+// CHECK: v1.5
+spv.module Logical GLSL450 requires #spv.vce<v1.5, [Shader], []> {
}
// -----
-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<v1.0, [Shader, Float16], []> {
+}
+
+// -----
+
+// CHECK: [SPV_KHR_float_controls, SPV_KHR_subgroup_vote]
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_float_controls, SPV_KHR_subgroup_vote]> {
}
diff --git a/mlir/test/Dialect/SPIRV/Serialization/non-uniform-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/non-uniform-ops.mlir
index cb3a40cc9c68..ab714dfbaa00 100644
--- 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<v1.0, [Shader], []> {
// 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
index 1435aaefb2ea..d4a46dd9f1f1 100644
--- 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<v1.0, [Shader], []> {
spv.func @foo() -> () "None" {
// CHECK: %[[CST:.*]] = spv.constant 0
%zero = spv.constant 0 : i32
@@ -17,15 +17,13 @@ spv.module "Logical" "GLSL450" {
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<v1.0, [Shader], []> {
spv.func @foo() -> () "None" {
// CHECK: %[[ZERO:.*]] = spv.constant 0
%zero = spv.constant 0 : i32
@@ -43,15 +41,13 @@ spv.module "Logical" "GLSL450" {
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<v1.0, [Shader], []> {
spv.func @foo() -> () "None" {
// CHECK: %[[CST0:.*]] = spv.constant 0
%zero = spv.constant 0 : i32
@@ -75,15 +71,13 @@ spv.module "Logical" "GLSL450" {
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<v1.0, [Shader], []> {
spv.func @foo() -> () "None" {
// CHECK: spv.Branch ^bb1
spv.Branch ^bb1
@@ -109,15 +103,13 @@ spv.module "Logical" "GLSL450" {
spv.Return
}
spv.EntryPoint "GLCompute" @main
-} attributes {
- capabilities = ["Shader"]
}
// -----
// Test multiple predecessors
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @foo() -> () "None" {
%var = spv.Variable : !spv.ptr<i32, Function>
@@ -160,15 +152,13 @@ spv.module "Logical" "GLSL450" {
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<v1.0, [Shader], []> {
spv.globalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
spv.globalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
spv.func @fmul_kernel() "None" {
@@ -245,4 +235,4 @@ spv.module "Logical" "GLSL450" {
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
index 6df3f0d4dd66..e391bae5b486 100644
--- 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<v1.0, [Shader], []> {
spv.func @selection(%cond: i1) -> () "None" {
// CHECK: spv.Branch ^bb1
// CHECK-NEXT: ^bb1:
@@ -48,8 +48,6 @@ spv.module "Logical" "GLSL450" {
}
spv.EntryPoint "GLCompute" @main
spv.ExecutionMode @main "LocalSize", 1, 1, 1
-} attributes {
- capabilities = ["Shader"]
}
// -----
@@ -57,7 +55,7 @@ spv.module "Logical" "GLSL450" {
// Selection with only then branch
// Selection in function entry block
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK: spv.func @selection(%[[ARG:.*]]: i1
spv.func @selection(%cond: i1) -> (i32) "None" {
// CHECK: spv.Branch ^bb1
@@ -87,7 +85,5 @@ spv.module "Logical" "GLSL450" {
}
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
index c905c68265cc..03cc85b8c087 100644
--- 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<v1.0, [Shader], []> {
// 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
index e477ffbd2f8a..e96cc418615f 100644
--- 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<v1.0, [Shader], []> {
// CHECK: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Input>
spv.globalVariable @var0 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Input>
diff --git a/mlir/test/Dialect/SPIRV/Serialization/terminator.mlir b/mlir/test/Dialect/SPIRV/Serialization/terminator.mlir
index 4127b0f9936a..e346d22f7af8 100644
--- 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<v1.0, [Shader], []> {
// 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
index 5f1f5b2a067f..6998930911db 100644
--- 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<v1.0, [Shader], []> {
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" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// 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
index a77fb11c4110..d8af9fa82607 100644
--- 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<vector<3xi32>, Input>
// CHECK-DAG: spv.globalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
@@ -122,4 +122,4 @@ spv.module "Logical" "GLSL450" {
}
// 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
index 1c1d698dc817..edc66c41591c 100644
--- 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<!spv.struct<f32 [0]>, StorageBuffer>
// CHECK-DAG: spv.globalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer>
// CHECK: spv.func [[FN:@.*]]()
@@ -24,4 +24,4 @@ spv.module "Logical" "GLSL450" {
}
// 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
index e0781defe25e..fc188c3938a8 100644
--- 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.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.module Logical GLSL450 {
spv.globalVariable @data bind(0, 0) : !spv.ptr<!spv.struct<!spv.rtarray<i32> [0]>, StorageBuffer>
spv.func @callee() "None" {
%0 = spv._address_of @data : !spv.ptr<!spv.struct<!spv.rtarray<i32> [0]>, StorageBuffer>
@@ -67,7 +67,7 @@ spv.module "Logical" "GLSL450" {
// -----
-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.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.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.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.module Logical GLSL450 {
spv.globalVariable @arg_0 bind(0, 0) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
spv.globalVariable @arg_1 bind(0, 1) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
@@ -222,7 +222,7 @@ spv.module "Logical" "GLSL450" {
}
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
index a73060fd327e..1129f89d7d84 100644
--- 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<!spv.struct<i32 [0], !spv.struct<f32 [0], i32 [4]> [4], f32 [12]>, Uniform>
spv.globalVariable @var0 bind(0,1) : !spv.ptr<!spv.struct<i32, !spv.struct<f32, i32>, f32>, Uniform>
@@ -31,7 +31,7 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
// CHECK: spv.globalVariable @var0 : !spv.ptr<!spv.struct<!spv.struct<!spv.struct<!spv.struct<!spv.struct<i1 [0], i1 [1], f64 [8]> [0], i1 [16]> [0], i1 [24]> [0], i1 [32]> [0], i1 [40]>, Uniform>
spv.globalVariable @var0 : !spv.ptr<!spv.struct<!spv.struct<!spv.struct<!spv.struct<!spv.struct<i1, i1, f64>, i1>, i1>, i1>, i1>, Uniform>
@@ -59,7 +59,7 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
// CHECK: spv.globalVariable @var0 : !spv.ptr<!spv.struct<vector<2xi32> [0], f32 [8]>, StorageBuffer>
spv.globalVariable @var0 : !spv.ptr<!spv.struct<vector<2xi32>, f32>, StorageBuffer>
@@ -72,7 +72,7 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
// CHECK: spv.globalVariable @emptyStructAsMember : !spv.ptr<!spv.struct<!spv.struct<> [0]>, StorageBuffer>
spv.globalVariable @emptyStructAsMember : !spv.ptr<!spv.struct<!spv.struct<>>, StorageBuffer>
@@ -91,7 +91,7 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
// CHECK: spv.globalVariable @var0 : !spv.ptr<!spv.struct<i32 [0]>, PushConstant>
spv.globalVariable @var0 : !spv.ptr<!spv.struct<i32>, PushConstant>
// CHECK: spv.globalVariable @var1 : !spv.ptr<!spv.struct<i32 [0]>, PhysicalStorageBuffer>
diff --git a/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir b/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir
index 4f43a77c48c9..60bf13e2571e 100644
--- 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<v1.0, [Shader], []>
-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<v1.0, [Shader], []>
+spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.5, [Shader], []>,
{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<v1.3, [GroupNonUniformBallot, Shader], []>
-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<v1.3, [GroupNonUniformBallot, Shader], []>
+spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.5, [Shader, GroupNonUniformBallot], []>,
{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 @@ spv.module "Logical" "GLSL450" {
// Test minimal capabilities.
-// CHECK: vce_triple = #spv.vce<v1.0, [Shader], []>
-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<v1.0, [Shader], []>
+spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader, Float16, Float64, Int16, Int64, VariablePointers], []>,
{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<v1.0, [Shader], []>
-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<v1.0, [Shader], []>
+spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [AtomicStorage], []>,
{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 @@ spv.module "Logical" "GLSL450" {
// * GroupNonUniformArithmetic
// * GroupNonUniformBallot
-// CHECK: vce_triple = #spv.vce<v1.3, [GroupNonUniformArithmetic, Shader], []>
-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<v1.3, [GroupNonUniformArithmetic, Shader], []>
+spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>,
{max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
-}
-
-// CHECK: vce_triple = #spv.vce<v1.3, [GroupNonUniformClustered, GroupNonUniformBallot, Shader], []>
-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<v1.3, [GroupNonUniformClustered, GroupNonUniformBallot, Shader], []>
+spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.3, [Shader, GroupNonUniformClustered, GroupNonUniformBallot], []>,
{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 @@ spv.module "Logical" "GLSL450" {
// Test deducing minimal extensions.
// spv.SubgroupBallotKHR requires the SPV_KHR_shader_ballot extension.
-// CHECK: vce_triple = #spv.vce<v1.0, [SubgroupBallotKHR, Shader], [SPV_KHR_shader_ballot]>
-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<v1.0, [SubgroupBallotKHR, Shader], [SPV_KHR_shader_ballot]>
+spv.module Logical GLSL450 attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader, SubgroupBallotKHR],
[SPV_KHR_shader_ballot, SPV_KHR_shader_clock, SPV_KHR_variable_pointers]>,
{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<v1.0, [VulkanMemoryModel], [SPV_KHR_vulkan_memory_model]>
-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<v1.0, [VulkanMemoryModel], [SPV_KHR_vulkan_memory_model]>
+spv.module Logical Vulkan attributes {
spv.target_env = #spv.target_env<
#spv.vce<v1.5, [Shader, VulkanMemoryModel], []>,
{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
index a5203a0e4a2a..e31c1bdeacca 100644
--- a/mlir/test/Dialect/SPIRV/availability.mlir
+++ b/mlir/test/Dialect/SPIRV/availability.mlir
@@ -36,7 +36,7 @@ func @module_logical_glsl450() {
// 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 @@ func @module_physical_storage_buffer64_vulkan() {
// 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
index 141d2c1aa1c3..9eab38cbcb05 100644
--- a/mlir/test/Dialect/SPIRV/control-flow-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/control-flow-ops.mlir
@@ -155,7 +155,7 @@ func @weights_cannot_both_be_zero() -> () {
// 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 @@ func @caller() {
// -----
-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.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.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.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.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.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 @@ func @in_other_func_like_op() {
// -----
// 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.module Logical GLSL450 {
spv.func @in_nested_region(%cond: i1) -> (i32) "None" {
spv.selection {
spv.BranchConditional %cond, ^then, ^merge
@@ -605,7 +605,7 @@ func @in_other_func_like_op(%arg: i32) -> i32 {
// -----
-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.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.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
index df8d2c80dcb0..aee4ff221dd9 100644
--- a/mlir/test/Dialect/SPIRV/ops.mlir
+++ b/mlir/test/Dialect/SPIRV/ops.mlir
@@ -416,7 +416,7 @@ func @u_convert_scalar(%arg0 : i32) -> i64 {
// spv.ExecutionMode
//===----------------------------------------------------------------------===//
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
@@ -425,7 +425,7 @@ spv.module "Logical" "GLSL450" {
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.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
@@ -639,7 +639,7 @@ func @aligned_load_incorrect_attributes() -> () {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.globalVariable @var0 : !spv.ptr<f32, Input>
// CHECK_LABEL: @simple_load
spv.func @simple_load() -> () "None" {
@@ -1057,7 +1057,7 @@ func @aligned_store_incorrect_attributes(%arg0 : f32) -> () {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.globalVariable @var0 : !spv.ptr<f32, Input>
spv.func @simple_store(%arg0 : f32) -> () "None" {
%0 = spv._address_of @var0 : !spv.ptr<f32, Input>
@@ -1130,7 +1130,7 @@ func @variable_init_normal_constant() -> () {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.globalVariable @global : !spv.ptr<f32, Workgroup>
spv.func @variable_init_global_variable() -> () "None" {
%0 = spv._address_of @global : !spv.ptr<f32, Workgroup>
@@ -1138,14 +1138,11 @@ spv.module "Logical" "GLSL450" {
%1 = spv.Variable init(%0) : !spv.ptr<!spv.ptr<f32, Workgroup>, 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
index 0247cac1b6fa..8a51ad5dc06d 100644
--- 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<!spv.struct<f32, !spv.array<4xf32>>, Input>
spv.func @access_chain() -> () "None" {
%0 = spv.constant 1: i32
@@ -28,7 +28,7 @@ func @address_of() -> () {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.globalVariable @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, 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.module Logical GLSL450 {
spv.globalVariable @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
spv.func @foo() -> () "None" {
// expected-error @+1 {{result type mismatch with the referenced global variable's type}}
@@ -135,7 +135,7 @@ func @value_result_num_elements_mismatch() -> () {
// spv.EntryPoint
//===----------------------------------------------------------------------===//
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
@@ -143,7 +143,7 @@ spv.module "Logical" "GLSL450" {
spv.EntryPoint "GLCompute" @do_nothing
}
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.globalVariable @var2 : !spv.ptr<f32, Input>
spv.globalVariable @var3 : !spv.ptr<f32, Output>
spv.func @do_something(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) -> () "None" {
@@ -157,7 +157,7 @@ spv.module "Logical" "GLSL450" {
// -----
-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.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
@@ -182,7 +182,7 @@ spv.module "Logical" "GLSL450" {
// -----
-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.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
@@ -202,12 +202,12 @@ spv.module "Logical" "GLSL450" {
// -----
-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 @@ spv.func @cannot_have_variadic_arguments(%arg: i32, ...) "None"
// -----
// 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.module "Logical" "GLSL450" {
// spv.globalVariable
//===----------------------------------------------------------------------===//
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
// CHECK: spv.globalVariable @var0 : !spv.ptr<f32, Input>
spv.globalVariable @var0 : !spv.ptr<f32, Input>
}
// 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<f32, Private>
// spv.globalVariable @var1 init(%0) : !spv.ptr<f32, Private>
@@ -280,7 +280,7 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.specConstant @sc = 4.0 : f32
// CHECK: spv.globalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
spv.globalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
@@ -295,13 +295,13 @@ spv.globalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
// CHECK: spv.globalVariable @var0 bind(1, 2) : !spv.ptr<f32, Uniform>
spv.globalVariable @var0 bind(1, 2) : !spv.ptr<f32, Uniform>
}
// 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<f32, Private>
// spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
@@ -309,7 +309,7 @@ spv.module "Logical" "GLSL450" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
// CHECK: spv.globalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
spv.globalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
// CHECK: spv.globalVariable @var2 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
@@ -326,28 +326,28 @@ module {
// -----
-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<f32, Private>
}
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
// expected-error @+1 {{storage class cannot be 'Generic'}}
spv.globalVariable @var0 : !spv.ptr<f32, Generic>
}
// -----
-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<f32, Input>
@@ -362,25 +362,33 @@ spv.module "Logical" "GLSL450" {
//===----------------------------------------------------------------------===//
// 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<v1.0, [Shader], [SPV_KHR_16bit_storage]>
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> { }
+
+// 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<v1.0, [Shader], [SPV_KHR_16bit_storage]> attributes {foo = "bar"}
+spv.module Logical GLSL450
+ requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]>
+ 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 @@ spv.module "Logical" "GLSL450" {
// -----
// Missing addressing model
-// expected-error at +1 {{custom op 'spv.module' expected addressing_model attribute specified as string}}
+// expected-error at +1 {{'spv.module' expected valid keyword}}
spv.module { }
// -----
// Wrong addressing model
-// expected-error at +1 {{custom op 'spv.module' invalid addressing_model attribute specification: "Physical"}}
-spv.module "Physical" { }
+// expected-error at +1 {{'spv.module' invalid addressing_model attribute specification: Physical}}
+spv.module Physical { }
// -----
// Missing memory model
-// expected-error at +1 {{custom op 'spv.module' expected memory_model attribute specified as string}}
-spv.module "Logical" { }
+// expected-error at +1 {{'spv.module' expected valid keyword}}
+spv.module Logical { }
// -----
// Wrong memory model
-// expected-error at +1 {{custom op 'spv.module' invalid memory_model attribute specification: "Bla"}}
-spv.module "Logical" "Bla" { }
+// expected-error at +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 @@ spv.module "Logical" "GLSL450" {
// -----
// 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 @@ spv.module "Logical" "GLSL450" {
// -----
// 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 @@ spv.module "Logical" "GLSL450" {
// -----
// 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 @@ func @module_end_not_in_module() -> () {
// 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 @@ func @reference_of() {
// -----
-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.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.module "Logical" "GLSL450" {
// 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" {
+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
index 32f36e96f5ea..9b42314e3f1d 100644
--- a/mlir/test/Dialect/SPIRV/target-env.mlir
+++ b/mlir/test/Dialect/SPIRV/target-env.mlir
@@ -148,7 +148,7 @@ func @subgroup_ballot_missing_extension(%predicate: i1) -> vector<4xi32> attribu
func @module_suitable_extension1() attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_EXT_physical_storage_buffer]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
- // CHECK: spv.module "PhysicalStorageBuffer64" "Vulkan"
+ // CHECK: spv.module PhysicalStorageBuffer64 Vulkan
"test.convert_to_module_op"() : () ->()
return
}
@@ -157,7 +157,7 @@ func @module_suitable_extension1() attributes {
func @module_suitable_extension2() attributes {
spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_KHR_physical_storage_buffer]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
- // CHECK: spv.module "PhysicalStorageBuffer64" "Vulkan"
+ // CHECK: spv.module PhysicalStorageBuffer64 Vulkan
"test.convert_to_module_op"() : () -> ()
return
}
@@ -185,7 +185,7 @@ 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<#spv.vce<v1.5, [VulkanMemoryModel, PhysicalStorageBufferAddresses], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
- // CHECK: spv.module "PhysicalStorageBuffer64" "Vulkan"
+ // 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
index 21f5c8cdd1e5..4ae375d63c55 100644
--- 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<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
+ {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
index 33f6472df4d2..f91bc71e8713 100644
--- a/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
+++ b/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
@@ -38,6 +38,7 @@ static LogicalResult runMLIRPasses(ModuleOp module) {
passManager.addPass(createConvertGPUToSPIRVPass());
OpPassManager &modulePM = passManager.nest<spirv::ModuleOp>();
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
index bee6a2d434b3..bfefebeefac3 100644
--- a/mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp
+++ b/mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp
@@ -63,7 +63,9 @@ class DeserializationTest : public ::testing::Test {
//===--------------------------------------------------------------------===//
/// 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<uint32_t> operands) {
diff --git a/mlir/unittests/Dialect/SPIRV/SerializationTest.cpp b/mlir/unittests/Dialect/SPIRV/SerializationTest.cpp
index f2831f123aec..c9f1e1570fa6 100644
--- 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 @@ class SerializationTest : public ::testing::Test {
state.addAttribute("memory_model",
builder.getI32IntegerAttr(
static_cast<uint32_t>(spirv::MemoryModel::GLSL450)));
+ state.addAttribute("vce_triple",
+ spirv::VerCapExtAttr::get(
+ spirv::Version::V_1_0, ArrayRef<spirv::Capability>(),
+ ArrayRef<spirv::Extension>(), &context));
spirv::ModuleOp::build(&builder, state);
module = cast<spirv::ModuleOp>(Operation::create(state));
}
More information about the Mlir-commits
mailing list