maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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
This commit is contained in:
Lei Zhang 2020-03-11 16:04:25 -04:00
parent c818c3cc96
commit 3148f10b17
53 changed files with 435 additions and 427 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td
View File

@ -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
//===----------------------------------------------------------------------===//

6
mlir/include/mlir/Dialect/SPIRV/SPIRVBinaryUtils.h
View File

@ -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);

1
mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h
View File

@ -23,6 +23,7 @@ namespace mlir {
class OpBuilder;
namespace spirv {
class VerCapExtAttr;
// TableGen'erated operation interfaces for querying versions, extensions, and
// capabilities.

41
mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td
View File

@ -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();
}

7
mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp
View File

@ -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 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,

171
mlir/lib/Dialect/SPIRV/SPIRVOps.cpp
View File

@ -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)) {

78
mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp
View File

@ -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 @@ private:
/// 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 @@ private:
/// 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 @@ private:
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) {

25
mlir/lib/Dialect/SPIRV/Serialization/SPIRVBinaryUtils.cpp
View File

@ -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)

29
mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp
View File

@ -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()))

2
mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp
View File

@ -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>>

16
mlir/test/Conversion/GPUToSPIRV/builtins.mlir
View File

@ -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()

2
mlir/test/Conversion/GPUToSPIRV/load-store.mlir
View File

@ -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>

3
mlir/test/Conversion/GPUToSPIRV/simple.mlir
View File

@ -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() {

4
mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
View File

@ -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

2
mlir/test/Dialect/SPIRV/Serialization/arithmetic-ops.mlir
View File

@ -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

4
mlir/test/Dialect/SPIRV/Serialization/array.mlir
View File

@ -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>

2
mlir/test/Dialect/SPIRV/Serialization/atomic-ops.mlir
View File

@ -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>

2
mlir/test/Dialect/SPIRV/Serialization/barrier.mlir
View File

@ -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"

2
mlir/test/Dialect/SPIRV/Serialization/bit-ops.mlir
View File

@ -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

4
mlir/test/Dialect/SPIRV/Serialization/cast-ops.mlir
View File

@ -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

2
mlir/test/Dialect/SPIRV/Serialization/composite-op.mlir
View File

@ -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>>

2
mlir/test/Dialect/SPIRV/Serialization/constant.mlir
View File

@ -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

4
mlir/test/Dialect/SPIRV/Serialization/entry-point.mlir
View File

@ -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"

2
mlir/test/Dialect/SPIRV/Serialization/execution-mode.mlir
View File

@ -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
}

2
mlir/test/Dialect/SPIRV/Serialization/function-call.mlir
View File

@ -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

6
mlir/test/Dialect/SPIRV/Serialization/global-variable.mlir
View File

@ -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>

2
mlir/test/Dialect/SPIRV/Serialization/glsl-ops.mlir
View File

@ -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

2
mlir/test/Dialect/SPIRV/Serialization/group-ops.mlir
View File

@ -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>

4
mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir
View File

@ -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

12
mlir/test/Dialect/SPIRV/Serialization/loop.mlir
View File

@ -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"]
}

6
mlir/test/Dialect/SPIRV/Serialization/memory-ops.mlir
View File

@ -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

24
mlir/test/Dialect/SPIRV/Serialization/module.mlir
View File

@ -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]> {
}

2
mlir/test/Dialect/SPIRV/Serialization/non-uniform-ops.mlir
View File

@ -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>

24
mlir/test/Dialect/SPIRV/Serialization/phi.mlir
View File

@ -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"]}
}

8
mlir/test/Dialect/SPIRV/Serialization/selection.mlir
View File

@ -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"]
}

2
mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir
View File

@ -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

2
mlir/test/Dialect/SPIRV/Serialization/struct.mlir
View File

@ -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>

2
mlir/test/Dialect/SPIRV/Serialization/terminator.mlir
View File

@ -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

4
mlir/test/Dialect/SPIRV/Serialization/undef.mlir
View File

@ -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

4
mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir
View File

@ -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"]}
}

4
mlir/test/Dialect/SPIRV/Transforms/abi-simple.mlir
View File

@ -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"]}
}

18
mlir/test/Dialect/SPIRV/Transforms/inlining.mlir
View File

@ -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.

10
mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir
View File

@ -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>

108
mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir
View File

@ -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
}
}

4
mlir/test/Dialect/SPIRV/availability.mlir
View File

@ -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
}

24
mlir/test/Dialect/SPIRV/control-flow-ops.mlir
View File

@ -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

17
mlir/test/Dialect/SPIRV/ops.mlir
View File

@ -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" {

124
mlir/test/Dialect/SPIRV/structure-ops.mlir
View File

@ -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@+1 {{custom op 'spv.module' expected addressing_model attribute specified as string}}
// expected-error@+1 {{'spv.module' expected valid keyword}}
spv.module { }
// -----
// Wrong addressing model
// expected-error@+1 {{custom op 'spv.module' invalid addressing_model attribute specification: "Physical"}}
spv.module "Physical" { }
// expected-error@+1 {{'spv.module' invalid addressing_model attribute specification: Physical}}
spv.module Physical { }
// -----
// Missing memory model
// expected-error@+1 {{custom op 'spv.module' expected memory_model attribute specified as string}}
spv.module "Logical" { }
// expected-error@+1 {{'spv.module' expected valid keyword}}
spv.module Logical { }
// -----
// Wrong memory model
// expected-error@+1 {{custom op 'spv.module' invalid memory_model attribute specification: "Bla"}}
spv.module "Logical" "Bla" { }
// expected-error@+1 {{'spv.module' invalid memory_model attribute specification: Bla}}
spv.module Logical Bla { }
// -----
// Module with multiple blocks
// expected-error @+1 {{expects region #0 to have 0 or 1 blocks}}
spv.module "Logical" "GLSL450" {
spv.module Logical GLSL450 {
^first:
spv.Return
^second:
@ -433,7 +441,7 @@ 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>
}

6
mlir/test/Dialect/SPIRV/target-env.mlir
View File

@ -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
}

8
mlir/test/mlir-vulkan-runner/addf.mlir
View File

@ -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>}} {

1
mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
View File

@ -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());

4
mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp
View File

@ -63,7 +63,9 @@ protected:
//===--------------------------------------------------------------------===//
/// 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) {

5
mlir/unittests/Dialect/SPIRV/SerializationTest.cpp
View File

@ -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 @@ protected:
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));
}