[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
//===----------------------------------------------------------------------===//
#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);
class OpBuilder;
namespace spirv {
+class VerCapExtAttr;
// TableGen'erated operation interfaces for querying versions, extensions, and
// capabilities.
### 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"]
}
```
}];
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.
let autogenSerialization = 0;
let extraClassDeclaration = [{
+ static StringRef getVCETripleAttrName() { return "vce_triple"; }
+
Block& getBlock() {
return this->getOperation()->getRegion(0).front();
}
PatternMatchResult GPUModuleConversion::matchAndRewrite(
gpu::GPUModuleOp moduleOp, ArrayRef<Value> operands,
ConversionPatternRewriter &rewriter) const {
- // TODO : Generalize this to account for different extensions,
- // capabilities, extended_instruction_sets, other addressing models
- // and memory models.
auto spvModule = rewriter.create<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,
#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"
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();
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(
}
spirv::MemoryAccess memoryAccessAttr;
- if (parseEnumAttribute(memoryAccessAttr, parser, state)) {
+ if (parseEnumStrAttr(memoryAccessAttr, parser, state)) {
return failure();
}
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();
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();
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();
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();
}
spirv::ExecutionMode execMode;
Attribute fn;
if (parser.parseAttribute(fn, kFnNameAttrName, state.attributes) ||
- parseEnumAttribute(execMode, parser, state)) {
+ parseEnumStrAttr(execMode, parser, state)) {
return failure();
}
// 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.
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)) {
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) {
// 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);
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");
}
}
- // 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();
}
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)) {
#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"
/// 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);
/// 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);
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;
}
}
- // Attaches the capabilities/extensions as an attribute to the module.
- attachCapabilities();
- attachExtensions();
+ attachVCETriple();
LLVM_DEBUG(llvm::dbgs() << "+++ completed deserialization +++\n");
return success();
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));
}
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();
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(
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();
}
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) {
//===----------------------------------------------------------------------===//
#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.
// | 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)
#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"
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());
}
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);
}
}
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()))
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>>
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()
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()
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()
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>}} {
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>}} {
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>}} {
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()
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()
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>
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{{[}][}]}}
// CHECK: spv.Return
gpu.return
}
- // CHECK: attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}
}
func @main() {
// 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>
}
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
// 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
// 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>
// -----
-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>
// 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>
// 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"
// 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
// 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
// -----
-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
// 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>>
// 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
// 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
}
// -----
-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"
// 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
}
// 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
// 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>
// -----
-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>
// -----
-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>
// 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
// 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>
// 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
// -----
-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
// 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
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" {
}
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) { }
// }
spv.Return
}
spv.EntryPoint "GLCompute" @main
-} attributes {
- capabilities = ["Shader"]
}
// 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
// -----
-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>
// -----
-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
// 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
}
// -----
-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]> {
}
// 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>
// 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
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
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
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
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>
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" {
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"]}
+}
// 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:
}
spv.EntryPoint "GLCompute" @main
spv.ExecutionMode @main "LocalSize", 1, 1, 1
-} attributes {
- capabilities = ["Shader"]
}
// -----
// 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
}
spv.EntryPoint "GLCompute" @main
spv.ExecutionMode @main "LocalSize", 1, 1, 1
-} attributes {
- capabilities = ["Shader"]
}
// 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
// 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>
// 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
// 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
// -----
-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
// 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")
}
// 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"]}
+}
// 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:@.*]]()
}
// CHECK: spv.EntryPoint "GLCompute" [[FN]]
// CHECK: spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
-} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}
+}
// 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
}
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @callee() -> i32 "None" {
%0 = spv.constant 42 : i32
spv.ReturnValue %0 : i32
// -----
-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>
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @callee(%cond : i1) -> () "None" {
spv.selection {
spv.BranchConditional %cond, ^then, ^merge
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @callee(%cond : i1) -> () "None" {
spv.selection {
spv.BranchConditional %cond, ^then, ^merge
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @callee(%cond : i1) -> () "None" {
spv.loop {
spv.Branch ^header
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @callee(%cond : i1) -> () "None" {
spv.loop {
spv.Branch ^header
// -----
-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>
}
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.
// 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>
// -----
-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>
// -----
-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>
// -----
-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>
// -----
-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>
// 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>
+ }
}
//===----------------------------------------------------------------------===//
// 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.
// * 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
+ }
}
//===----------------------------------------------------------------------===//
// 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
+ }
}
// 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
}
// 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
}
// 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>
// -----
-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)
// -----
-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)
// -----
-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) -> ()
// -----
-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}}
// -----
-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'}}
// -----
-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
// -----
// 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
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @in_nested_region(%cond: i1) -> (i32) "None" {
spv.selection {
spv.BranchConditional %cond, ^then, ^merge
// -----
-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}}
// -----
-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')}}
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @in_nested_region(%cond: i1) -> () "None" {
spv.selection {
spv.BranchConditional %cond, ^then, ^merge
// spv.ExecutionMode
//===----------------------------------------------------------------------===//
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
spv.ExecutionMode @do_nothing "ContractionOff"
}
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.globalVariable @var0 : !spv.ptr<f32, Input>
// CHECK_LABEL: @simple_load
spv.func @simple_load() -> () "None" {
// -----
-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>
// -----
-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>
%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" {
// 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
// -----
-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}}
// -----
-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}}
// spv.EntryPoint
//===----------------------------------------------------------------------===//
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
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" {
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
// -----
-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
// -----
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
spv.func @do_nothing() -> () "None" {
spv.Return
}
// -----
-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
}
// -----
// 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" {
// 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>
// -----
-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>
// -----
-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>
// -----
-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>
// -----
-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>
//===----------------------------------------------------------------------===//
// 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
}
// -----
// Missing addressing model
-// expected-error@+1 {{custom op 'spv.module' expected addressing_model attribute specified as string}}
+// expected-error@+1 {{'spv.module' expected valid keyword}}
spv.module { }
// -----
// Wrong addressing model
-// expected-error@+1 {{custom op 'spv.module' invalid addressing_model attribute specification: "Physical"}}
-spv.module "Physical" { }
+// expected-error@+1 {{'spv.module' invalid addressing_model attribute specification: Physical}}
+spv.module Physical { }
// -----
// Missing memory model
-// expected-error@+1 {{custom op 'spv.module' expected memory_model attribute specified as string}}
-spv.module "Logical" { }
+// expected-error@+1 {{'spv.module' expected valid keyword}}
+spv.module Logical { }
// -----
// Wrong memory model
-// expected-error@+1 {{custom op 'spv.module' invalid memory_model attribute specification: "Bla"}}
-spv.module "Logical" "Bla" { }
+// expected-error@+1 {{'spv.module' invalid memory_model attribute specification: Bla}}
+spv.module Logical Bla { }
// -----
// Module with multiple blocks
// expected-error @+1 {{expects region #0 to have 0 or 1 blocks}}
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
^first:
spv.Return
^second:
// -----
// 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"() : () -> ()
}
// -----
// 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"() : () -> ()
// -----
// 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
//===----------------------------------------------------------------------===//
// 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
// -----
-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
// -----
-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}}
// 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
// -----
-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>
}
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
}
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
}
// 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
}
// 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>}} {
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());
//===--------------------------------------------------------------------===//
/// 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) {
//===----------------------------------------------------------------------===//
#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"
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));
}
projects / platform / upstream / llvm.git / commitdiff