Update MLIR API.
PiperOrigin-RevId: 300753498
diff --git a/iree/compiler/Dialect/Flow/Analysis/Dispatchability.cpp b/iree/compiler/Dialect/Flow/Analysis/Dispatchability.cpp
index 1d2951a..808382e 100644
--- a/iree/compiler/Dialect/Flow/Analysis/Dispatchability.cpp
+++ b/iree/compiler/Dialect/Flow/Analysis/Dispatchability.cpp
@@ -115,10 +115,6 @@
// TODO(benvanik): widen to all known terminators? sometimes they may
// have side-effects.
continue;
- } else if (!op.getDialect() || !op.hasNoSideEffect()) {
- // Ops with side-effects cannot be dispatched as we must be able to
- // exactly model I/O.
- return false;
} else if (isa<xla_hlo::DotOp>(op) || isa<xla_hlo::ConvOp>(op)) {
// Some unfusable ops must remain on their own.
return false;
@@ -126,6 +122,13 @@
isa<xla_hlo::ReduceWindowOp>(op)) {
// Reductions always become flow ops.
return false;
+
+ // TODO: Properly handle region side effects.
+ } else if (!MemoryEffectOpInterface::hasNoEffect(&op) ||
+ op.getNumRegions() != 0) {
+ // Ops with side-effects cannot be dispatched as we must be able to
+ // exactly model I/O.
+ return false;
}
}
}
diff --git a/iree/compiler/Dialect/Flow/IR/BUILD b/iree/compiler/Dialect/Flow/IR/BUILD
index 2515f41..eace239 100644
--- a/iree/compiler/Dialect/Flow/IR/BUILD
+++ b/iree/compiler/Dialect/Flow/IR/BUILD
@@ -100,5 +100,6 @@
":td_files",
"//iree/compiler/Dialect/IREE/IR:td_files",
"@llvm-project//mlir:OpBaseTdFiles",
+ "@llvm-project//mlir:include/mlir/Interfaces/SideEffects.td",
],
)
diff --git a/iree/compiler/Dialect/Flow/IR/FlowBase.td b/iree/compiler/Dialect/Flow/IR/FlowBase.td
index 725675a..80cb26c 100644
--- a/iree/compiler/Dialect/Flow/IR/FlowBase.td
+++ b/iree/compiler/Dialect/Flow/IR/FlowBase.td
@@ -78,9 +78,6 @@
let printer = [{ return print$cppClass(p, *this); }];
}
-class FLOW_PureOp<string mnemonic, list<OpTrait> traits = []> :
- FLOW_Op<mnemonic, !listconcat(traits, [NoSideEffect])>;
-
def FLOW_StreamableOp : OpInterface<"StreamableOpInterface"> {
let description = [{
Interface for ops that can be used within a stream.
diff --git a/iree/compiler/Dialect/Flow/IR/FlowOps.td b/iree/compiler/Dialect/Flow/IR/FlowOps.td
index 23a5616..8cc2335 100644
--- a/iree/compiler/Dialect/Flow/IR/FlowOps.td
+++ b/iree/compiler/Dialect/Flow/IR/FlowOps.td
@@ -16,6 +16,10 @@
#define IREE_DIALECT_FLOW_OPS
include "iree/compiler/Dialect/Flow/IR/FlowBase.td"
+include "mlir/Interfaces/SideEffects.td"
+
+class FLOW_PureOp<string mnemonic, list<OpTrait> traits = []> :
+ FLOW_Op<mnemonic, !listconcat(traits, [NoSideEffect])>;
//===----------------------------------------------------------------------===//
// Variables
diff --git a/iree/compiler/Dialect/Flow/Transforms/FormStreams.cpp b/iree/compiler/Dialect/Flow/Transforms/FormStreams.cpp
index f7d75b4..e4f21de 100644
--- a/iree/compiler/Dialect/Flow/Transforms/FormStreams.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/FormStreams.cpp
@@ -68,7 +68,9 @@
// we need to bail.
SmallVector<Operation *, 8> opsWithSideEffects;
for (auto &op : block) {
- if (!op.hasNoSideEffect() || op.hasTrait<OpTrait::IREE::YieldPoint>()) {
+ // TODO: Handle region side effects.
+ if (!MemoryEffectOpInterface::hasNoEffect(&op) ||
+ op.hasTrait<OpTrait::IREE::YieldPoint>() || op.getNumRegions() != 0) {
opsWithSideEffects.push_back(&op);
}
}
@@ -111,7 +113,9 @@
}
processedOps.insert(op);
- if (!op->hasNoSideEffect()) {
+ // TODO: Handle region side effects.
+ if (!MemoryEffectOpInterface::hasNoEffect(op) ||
+ op->getNumRegions() != 0) {
// Op has side-effects and should split the current stream.
resetCurrentStream();
return;
@@ -140,7 +144,10 @@
// cross-block dependencies.
markOpDAGOutside(depOp);
continue;
- } else if (!depOp->hasNoSideEffect() || !isStreamableOp(depOp)) {
+
+ // TODO: Handle region side effects.
+ } else if (!MemoryEffectOpInterface::hasNoEffect(depOp) ||
+ !isStreamableOp(depOp) || depOp->getNumRegions() != 0) {
// Source op has side effects or isn't streamable meaning that we
// can't fold it into the stream region.
markOpDAGOutside(depOp);
diff --git a/iree/compiler/Dialect/HAL/IR/HALBase.td b/iree/compiler/Dialect/HAL/IR/HALBase.td
index 5294d88..5660b5b 100644
--- a/iree/compiler/Dialect/HAL/IR/HALBase.td
+++ b/iree/compiler/Dialect/HAL/IR/HALBase.td
@@ -563,10 +563,4 @@
let printer = [{ return print$cppClass(p, *this); }];
}
-class HAL_PureOp<string mnemonic, list<OpTrait> traits = []> :
- HAL_Op<mnemonic, !listconcat(traits, [NoSideEffect])>;
-
-class HAL_MakeTupleOp<string mnemonic, list<OpTrait> traits = []> :
- HAL_PureOp<mnemonic, traits>;
-
#endif // IREE_DIALECT_HAL_BASE
diff --git a/iree/compiler/Dialect/HAL/IR/HALOps.td b/iree/compiler/Dialect/HAL/IR/HALOps.td
index ed0df41..a900ddd 100644
--- a/iree/compiler/Dialect/HAL/IR/HALOps.td
+++ b/iree/compiler/Dialect/HAL/IR/HALOps.td
@@ -17,6 +17,13 @@
include "iree/compiler/Dialect/HAL/IR/HALBase.td"
include "mlir/IR/OpAsmInterface.td"
+include "mlir/Interfaces/SideEffects.td"
+
+class HAL_PureOp<string mnemonic, list<OpTrait> traits = []> :
+ HAL_Op<mnemonic, !listconcat(traits, [NoSideEffect])>;
+
+class HAL_MakeTupleOp<string mnemonic, list<OpTrait> traits = []> :
+ HAL_PureOp<mnemonic, traits>;
//===----------------------------------------------------------------------===//
// Magic temporary hacks
diff --git a/iree/compiler/Dialect/HAL/Target/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/test/smoketest.mlir
index d9b2169..55ddce3 100644
--- a/iree/compiler/Dialect/HAL/Target/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/test/smoketest.mlir
@@ -21,15 +21,13 @@
// VKSPV-NEXT: hal.executable.binary attributes {
// VKSPV-SAME: data = dense
// VKSPV-SAME: format = 1397773893 : i32} {
-// VKSPV-NEXT: module attributes {spv.target_env = #spv.target_env<V_1_3, [SPV_KHR_storage_buffer_storage_class], [Shader], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
-// VKSPV-NEXT: spv.module "Logical" "GLSL450" {
+// VKSPV-NEXT: module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
+// VKSPV-NEXT: spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
// VKSPV-DAG: spv.globalVariable [[GLOBALID:@.*]] built_in("GlobalInvocationId")
// VKSPV-DAG: spv.globalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
// VKSPV: spv.EntryPoint "GLCompute" @simpleMath_rgn_dispatch_0, [[GLOBALID]], [[NUMWORKGROUPS]]
// VKSPV-NEXT: spv.ExecutionMode @simpleMath_rgn_dispatch_0 "LocalSize", 32, 1, 1
-// VKSPV-NEXT: } attributes {
-// VKSPV-SAME: capabilities = ["Shader"],
-// VKSPV-SAME: extensions = ["SPV_KHR_storage_buffer_storage_class"]}
+// VKSPV-NEXT: }
// VKSPV-NEXT: }
// VKSPV-NEXT: }
// VKSPV-NEXT: }
diff --git a/iree/compiler/Dialect/IREE/IR/BUILD b/iree/compiler/Dialect/IREE/IR/BUILD
index 2e32b25..d292976 100644
--- a/iree/compiler/Dialect/IREE/IR/BUILD
+++ b/iree/compiler/Dialect/IREE/IR/BUILD
@@ -65,5 +65,6 @@
td_srcs = [
":td_files",
"@llvm-project//mlir:OpBaseTdFiles",
+ "@llvm-project//mlir:include/mlir/Interfaces/SideEffects.td",
],
)
diff --git a/iree/compiler/Dialect/IREE/IR/IREEOps.td b/iree/compiler/Dialect/IREE/IR/IREEOps.td
index 2acc2cb..d200a80 100644
--- a/iree/compiler/Dialect/IREE/IR/IREEOps.td
+++ b/iree/compiler/Dialect/IREE/IR/IREEOps.td
@@ -16,6 +16,7 @@
#define IREE_DIALECT_IREE_OPS
include "iree/compiler/Dialect/IREE/IR/IREEBase.td"
+include "mlir/Interfaces/SideEffects.td"
//===----------------------------------------------------------------------===//
// Op types
diff --git a/iree/compiler/Dialect/Modules/Strings/IR/Ops.td b/iree/compiler/Dialect/Modules/Strings/IR/Ops.td
index bfcd195..7846c30 100644
--- a/iree/compiler/Dialect/Modules/Strings/IR/Ops.td
+++ b/iree/compiler/Dialect/Modules/Strings/IR/Ops.td
@@ -17,6 +17,7 @@
include "iree/compiler/Dialect/IREE/IR/IREEBase.td"
include "mlir/IR/OpBase.td"
+include "mlir/Interfaces/SideEffects.td"
def IREESTRINGS_Dialect : Dialect {
let name = "strings";
diff --git a/iree/compiler/Dialect/Shape/IR/BUILD b/iree/compiler/Dialect/Shape/IR/BUILD
index fa0bbcc..2952c4d 100644
--- a/iree/compiler/Dialect/Shape/IR/BUILD
+++ b/iree/compiler/Dialect/Shape/IR/BUILD
@@ -69,5 +69,6 @@
":td_files",
"//iree/compiler/Dialect/IREE/IR:td_files",
"@llvm-project//mlir:OpBaseTdFiles",
+ "@llvm-project//mlir:include/mlir/Interfaces/SideEffects.td",
],
)
diff --git a/iree/compiler/Dialect/Shape/IR/ShapeOps.td b/iree/compiler/Dialect/Shape/IR/ShapeOps.td
index f6083ab..2092d83 100644
--- a/iree/compiler/Dialect/Shape/IR/ShapeOps.td
+++ b/iree/compiler/Dialect/Shape/IR/ShapeOps.td
@@ -16,6 +16,7 @@
#define IREE_DIALECT_SHAPE_OPS
include "iree/compiler/Dialect/Shape/IR/ShapeBase.td"
+include "mlir/Interfaces/SideEffects.td"
//===----------------------------------------------------------------------===//
// Op types
diff --git a/iree/compiler/Dialect/Shape/Transforms/HoistShapeCalculations.cpp b/iree/compiler/Dialect/Shape/Transforms/HoistShapeCalculations.cpp
index 2f78ee2..9d06d6b 100644
--- a/iree/compiler/Dialect/Shape/Transforms/HoistShapeCalculations.cpp
+++ b/iree/compiler/Dialect/Shape/Transforms/HoistShapeCalculations.cpp
@@ -29,7 +29,7 @@
bool isSimpleShapeCalculationOp(Operation *op) {
// The op must have no side effects.
- if (!op->hasNoSideEffect()) {
+ if (!MemoryEffectOpInterface::hasNoEffect(op) || op->getNumRegions() != 0) {
return false;
}
// The op should operate on types that are likely shape calculations.
diff --git a/iree/compiler/Dialect/VM/IR/BUILD b/iree/compiler/Dialect/VM/IR/BUILD
index 75c7391..68253d0 100644
--- a/iree/compiler/Dialect/VM/IR/BUILD
+++ b/iree/compiler/Dialect/VM/IR/BUILD
@@ -78,8 +78,6 @@
":td_files",
"//iree/compiler/Dialect/IREE/IR:td_files",
"@llvm-project//mlir:OpBaseTdFiles",
- "@llvm-project//mlir:include/mlir/Interfaces/CallInterfaces.td",
- "@llvm-project//mlir:include/mlir/Interfaces/ControlFlowInterfaces.td",
],
)
@@ -97,6 +95,7 @@
"@llvm-project//mlir:OpBaseTdFiles",
"@llvm-project//mlir:include/mlir/Interfaces/CallInterfaces.td",
"@llvm-project//mlir:include/mlir/Interfaces/ControlFlowInterfaces.td",
+ "@llvm-project//mlir:include/mlir/Interfaces/SideEffects.td",
],
)
@@ -113,6 +112,7 @@
"@llvm-project//mlir:OpBaseTdFiles",
"@llvm-project//mlir:include/mlir/Interfaces/CallInterfaces.td",
"@llvm-project//mlir:include/mlir/Interfaces/ControlFlowInterfaces.td",
+ "@llvm-project//mlir:include/mlir/Interfaces/SideEffects.td",
],
)
@@ -128,7 +128,5 @@
":td_files",
"//iree/compiler/Dialect/IREE/IR:td_files",
"@llvm-project//mlir:OpBaseTdFiles",
- "@llvm-project//mlir:include/mlir/Interfaces/CallInterfaces.td",
- "@llvm-project//mlir:include/mlir/Interfaces/ControlFlowInterfaces.td",
],
)
diff --git a/iree/compiler/Dialect/VM/IR/VMBase.td b/iree/compiler/Dialect/VM/IR/VMBase.td
index ed33a50..3e6c4ee 100644
--- a/iree/compiler/Dialect/VM/IR/VMBase.td
+++ b/iree/compiler/Dialect/VM/IR/VMBase.td
@@ -343,9 +343,6 @@
list<VM_EncEncodeExpr> encoding = ?;
}
-class VM_PureOp<string mnemonic, list<OpTrait> traits = []> :
- VM_Op<mnemonic, !listconcat(traits, [NoSideEffect])>;
-
//===----------------------------------------------------------------------===//
// VM traits
//===----------------------------------------------------------------------===//
diff --git a/iree/compiler/Dialect/VM/IR/VMOps.td b/iree/compiler/Dialect/VM/IR/VMOps.td
index 9d821f2..a2143d0 100644
--- a/iree/compiler/Dialect/VM/IR/VMOps.td
+++ b/iree/compiler/Dialect/VM/IR/VMOps.td
@@ -17,8 +17,12 @@
include "mlir/Interfaces/CallInterfaces.td"
include "mlir/Interfaces/ControlFlowInterfaces.td"
+include "mlir/Interfaces/SideEffects.td"
include "iree/compiler/Dialect/VM/IR/VMBase.td"
+class VM_PureOp<string mnemonic, list<OpTrait> traits = []> :
+ VM_Op<mnemonic, !listconcat(traits, [NoSideEffect])>;
+
//===----------------------------------------------------------------------===//
// Structural ops
//===----------------------------------------------------------------------===//
@@ -521,6 +525,7 @@
class VM_ConstOp<string mnemonic, string ctype, list<OpTrait> traits = []> :
VM_PureOp<mnemonic, !listconcat(traits, [
+ ConstantLike,
DeclareOpInterfaceMethods<VM_SerializableOpInterface>,
])> {
let skipDefaultBuilders = 1;
@@ -571,6 +576,7 @@
}
def VM_ConstI32ZeroOp : VM_PureOp<"const.i32.zero", [
+ ConstantLike,
DeclareOpInterfaceMethods<VM_SerializableOpInterface>,
]> {
let summary = [{32-bit integer constant zero operation}];
@@ -600,6 +606,7 @@
}
def VM_ConstRefZeroOp : VM_PureOp<"const.ref.zero", [
+ ConstantLike,
DeclareOpInterfaceMethods<VM_SerializableOpInterface>,
]> {
let summary = [{null ref_ptr constant operation}];
diff --git a/iree/compiler/Dialect/VMLA/IR/BUILD b/iree/compiler/Dialect/VMLA/IR/BUILD
index 353ac92..a4be0d3 100644
--- a/iree/compiler/Dialect/VMLA/IR/BUILD
+++ b/iree/compiler/Dialect/VMLA/IR/BUILD
@@ -94,7 +94,6 @@
"//iree/compiler/Dialect/IREE/IR:td_files",
"//iree/compiler/Dialect/Shape/IR:td_files",
"@llvm-project//mlir:OpBaseTdFiles",
- "@llvm-project//mlir:include/mlir/Interfaces/CallInterfaces.td",
"@llvm-project//mlir:StdOpsTdFiles",
],
)
@@ -112,7 +111,6 @@
"//iree/compiler/Dialect/IREE/IR:td_files",
"//iree/compiler/Dialect/Shape/IR:td_files",
"@llvm-project//mlir:OpBaseTdFiles",
- "@llvm-project//mlir:include/mlir/Interfaces/CallInterfaces.td",
"@llvm-project//mlir:StdOpsTdFiles",
],
)
@@ -130,6 +128,5 @@
"//iree/compiler/Dialect/IREE/IR:td_files",
"//iree/compiler/Dialect/Shape/IR:td_files",
"@llvm-project//mlir:OpBaseTdFiles",
- "@llvm-project//mlir:include/mlir/Interfaces/CallInterfaces.td",
],
)
diff --git a/iree/compiler/Dialect/VMLA/IR/VMLABase.td b/iree/compiler/Dialect/VMLA/IR/VMLABase.td
index 799e4af..ae0d561 100644
--- a/iree/compiler/Dialect/VMLA/IR/VMLABase.td
+++ b/iree/compiler/Dialect/VMLA/IR/VMLABase.td
@@ -151,9 +151,6 @@
// let printer = [{ return print$cppClass(p, *this); }];
}
-class VMLA_PureOp<string mnemonic, list<OpTrait> traits = []> :
- VMLA_Op<mnemonic, !listconcat(traits, [NoSideEffect])>;
-
class VMLA_ElementTypeOp<string mnemonic, list<OpTrait> traits = []> :
VMLA_Op<mnemonic, !listconcat(traits, [VMLA_OpInterface])> {
let extraClassDeclaration = [{
diff --git a/iree/compiler/Dialect/VMLA/IR/VMLAOps.td b/iree/compiler/Dialect/VMLA/IR/VMLAOps.td
index 5e796bd..c82e660 100644
--- a/iree/compiler/Dialect/VMLA/IR/VMLAOps.td
+++ b/iree/compiler/Dialect/VMLA/IR/VMLAOps.td
@@ -17,6 +17,10 @@
include "iree/compiler/Dialect/VMLA/IR/VMLABase.td"
include "mlir/IR/OpAsmInterface.td"
+include "mlir/Interfaces/SideEffects.td"
+
+class VMLA_PureOp<string mnemonic, list<OpTrait> traits = []> :
+ VMLA_Op<mnemonic, !listconcat(traits, [NoSideEffect])>;
//===----------------------------------------------------------------------===//
// VMLA Ops: pseudo ops
diff --git a/iree/compiler/Dialect/Vulkan/Utils/TargetEnvUtils.cpp b/iree/compiler/Dialect/Vulkan/Utils/TargetEnvUtils.cpp
index 139c9b7..fd3b6ea 100644
--- a/iree/compiler/Dialect/Vulkan/Utils/TargetEnvUtils.cpp
+++ b/iree/compiler/Dialect/Vulkan/Utils/TargetEnvUtils.cpp
@@ -103,29 +103,6 @@
vkCapabilities.maxComputeWorkGroupInvocations(),
vkCapabilities.maxComputeWorkGroupSize(), context);
}
-
-// TODO(antiagainst): move this into MLIR core as spirv::TargetEnvAttr::get().
-spirv::TargetEnvAttr getSpirvTargetEnv(spirv::Version version,
- ArrayRef<spirv::Extension> extensions,
- ArrayRef<spirv::Capability> capabilities,
- DictionaryAttr limits) {
- Builder b(limits.getContext());
-
- auto versionAttr = b.getI32IntegerAttr(static_cast<uint32_t>(version));
-
- SmallVector<Attribute, 4> extAttrs;
- extAttrs.reserve(extensions.size());
- for (spirv::Extension ext : extensions)
- extAttrs.push_back(b.getStringAttr(spirv::stringifyExtension(ext)));
-
- SmallVector<Attribute, 4> capAttrs;
- capAttrs.reserve(capabilities.size());
- for (spirv::Capability cap : capabilities)
- capAttrs.push_back(b.getI32IntegerAttr(static_cast<uint32_t>(cap)));
-
- return spirv::TargetEnvAttr::get(versionAttr, b.getArrayAttr(extAttrs),
- b.getArrayAttr(capAttrs), limits);
-}
} // anonymous namespace
// TODO(antiagainst): register more SwiftShader extensions.
@@ -146,8 +123,9 @@
auto spvLimits = convertResourceLimits(vkTargetEnv);
- return getSpirvTargetEnv(spvVersion, spvExtensions, spvCapabilities,
- spvLimits);
+ auto triple = spirv::VerCapExtAttr::get(
+ spvVersion, spvCapabilities, spvExtensions, vkTargetEnv.getContext());
+ return spirv::TargetEnvAttr::get(triple, spvLimits);
}
} // namespace Vulkan
diff --git a/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir b/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir
index c3a8ae2..5867d2c 100644
--- a/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir
+++ b/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir
@@ -1,12 +1,12 @@
// RUN: iree-opt -iree-hal-transformation-pipeline -iree-hal-target-backends=vulkan-spirv %s | IreeFileCheck %s -check-prefix=DEFAULT
-// RUN: iree-opt -iree-hal-transformation-pipeline -iree-hal-target-backends=vulkan-spirv -iree-vulkan-target-env="#vk.target_env<v1.0, r(10), [VK_KHR_shader_float16_int8], {maxComputeWorkGroupInvocations = 64: i32, maxComputeWorkGroupSize = dense<[8, 8, 8]>: vector<3xi32>, shaderFloat64, shaderInt8}>" %s | IreeFileCheck %s -check-prefix=V10
+// RUN: iree-opt -iree-hal-transformation-pipeline -iree-hal-target-backends=vulkan-spirv -iree-vulkan-target-env="#vk.target_env<v1.0, r(10), [VK_KHR_shader_float16_int8, VK_KHR_storage_buffer_storage_class], {maxComputeWorkGroupInvocations = 64: i32, maxComputeWorkGroupSize = dense<[8, 8, 8]>: vector<3xi32>, shaderFloat64, shaderInt8}>" %s | IreeFileCheck %s -check-prefix=V10
// TODO(antiagainst): Passing in lenghty strings as command-line options is not
// optimal. We should consider creating a dedicated test pass to pick up
// #vk.target_env in input assembly and convert them.
-// DEFAULT: spv.target_env = #spv.target_env<V_1_3, [SPV_KHR_storage_buffer_storage_class], [Shader], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>}
-// V10: spv.target_env = #spv.target_env<V_1_0, [], [Shader, Float64, Int8], {max_compute_workgroup_invocations = 64 : i32, max_compute_workgroup_size = dense<8> : vector<3xi32>}>
+// DEFAULT: spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>}
+// V10: spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader, Float64, Int8], [SPV_KHR_storage_buffer_storage_class]>, {max_compute_workgroup_invocations = 64 : i32, max_compute_workgroup_size = dense<8> : vector<3xi32>}>
flow.executable @simpleMath_ex_dispatch_0 {
flow.dispatch.entry @simpleMath_rgn_dispatch_0 attributes {
workload = dense<[4, 1, 1]> : vector<3xi32>
diff --git a/iree/compiler/Translation/SPIRV/LinalgToSPIRV/LowerToSPIRV.cpp b/iree/compiler/Translation/SPIRV/LinalgToSPIRV/LowerToSPIRV.cpp
index f79d227..4d1178e 100644
--- a/iree/compiler/Translation/SPIRV/LinalgToSPIRV/LowerToSPIRV.cpp
+++ b/iree/compiler/Translation/SPIRV/LinalgToSPIRV/LowerToSPIRV.cpp
@@ -338,6 +338,7 @@
spirvModulePM.addPass(spirv::createLowerABIAttributesPass());
spirvModulePM.addPass(createCanonicalizerPass());
spirvModulePM.addPass(createCSEPass());
+ spirvModulePM.addPass(spirv::createUpdateVersionCapabilityExtensionPass());
}
void addLowerToSPIRVPasses(OpPassManager &pm, ArrayRef<int64_t> workGroupSize) {
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/IREEToSPIRVPass.cpp b/iree/compiler/Translation/SPIRV/XLAToSPIRV/IREEToSPIRVPass.cpp
index 17fe59a..3779b27 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/IREEToSPIRVPass.cpp
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/IREEToSPIRVPass.cpp
@@ -176,8 +176,7 @@
// Create a spirv.module Op.
auto spvModule = builder.create<spirv::ModuleOp>(
module.getLoc(), spirv::AddressingModel::Logical,
- spirv::MemoryModel::GLSL450, spirv::Capability::Shader,
- spirv::Extension::SPV_KHR_storage_buffer_storage_class);
+ spirv::MemoryModel::GLSL450);
// Generate the SPIR-V entry function for the dispatch function
if (failed(lowerEntryFunctions(spvModule, fns))) {
@@ -209,6 +208,7 @@
OpPassManager &spirvPasses = conversionPassManager.nest<spirv::ModuleOp>();
spirvPasses.addPass(spirv::createLowerABIAttributesPass());
spirvPasses.addPass(createAdjustIntegerWidthPass());
+ spirvPasses.addPass(spirv::createUpdateVersionCapabilityExtensionPass());
}
} // namespace iree_compiler
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/adjust_integer_width.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/adjust_integer_width.mlir
index b9c2ef5..619e42d 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/adjust_integer_width.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/adjust_integer_width.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -iree-spirv-adjust-integer-width -verify-diagnostics -o - %s | IreeFileCheck %s
module{
- spv.module "Logical" "GLSL450" {
+ spv.module Logical GLSL450 {
spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
// CHECK: spv.globalVariable @constant_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
// CHECK: spv.globalVariable @constant_arg_1 bind(0, 1) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
@@ -26,7 +26,7 @@
}
}
- spv.module "Logical" "GLSL450" {
+ spv.module Logical GLSL450 {
spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
// CHECK: spv.globalVariable @constant_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
// CHECK: spv.globalVariable @constant_arg_1 bind(0, 1) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
@@ -50,7 +50,7 @@
}
}
- spv.module "Logical" "GLSL450" {
+ spv.module Logical GLSL450 {
spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
// CHECK: spv.globalVariable @constant_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
// CHECK: spv.globalVariable @constant_arg_1 bind(0, 1) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
@@ -74,7 +74,7 @@
}
}
- spv.module "Logical" "GLSL450" {
+ spv.module Logical GLSL450 {
spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
// CHECK: spv.globalVariable @constant_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
// CHECK: spv.globalVariable @constant_arg_1 bind(0, 1) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
@@ -101,7 +101,7 @@
}
}
- spv.module "Logical" "GLSL450" {
+ spv.module Logical GLSL450 {
spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
spv.globalVariable @arg_0 bind(0, 0) : !spv.ptr<!spv.struct<i64 [0]>, StorageBuffer>
spv.globalVariable @arg_1 bind(0, 1) : !spv.ptr<!spv.struct<i64 [0]>, StorageBuffer>
@@ -273,7 +273,7 @@
}
}
- spv.module "Logical" "GLSL450" {
+ spv.module Logical GLSL450 {
spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
spv.globalVariable @constant_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<i1 [0]>, StorageBuffer>
spv.globalVariable @constant_arg_1 bind(0, 1) : !spv.ptr<!spv.struct<i8 [0]>, StorageBuffer>
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast.mlir
index da10900..3186c6a 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -iree-index-computation -simplify-spirv-affine-exprs=false -convert-iree-to-spirv -verify-diagnostics -o - %s | IreeFileCheck %s
module {
- // CHECK:spv.module "Logical" "GLSL450"
+ // CHECK:spv.module Logical GLSL450
// CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
// CHECK: spv.func [[FN:@broadcast_2D_3D]]
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
@@ -23,7 +23,7 @@
// -----
module {
- // CHECK:spv.module "Logical" "GLSL450"
+ // CHECK:spv.module Logical GLSL450
// CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
// CHECK: spv.func [[FN:@broadcast_scalar_3D]]
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast_in_dim.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast_in_dim.mlir
index 52f9bdf..0dd1965 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast_in_dim.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast_in_dim.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -iree-index-computation -simplify-spirv-affine-exprs=false -convert-iree-to-spirv -verify-diagnostics -o - %s | IreeFileCheck %s
module {
- // CHECK:spv.module "Logical" "GLSL450"
+ // CHECK:spv.module Logical GLSL450
// CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
// CHECK: spv.func [[FN:@broadcast_in_dim_2D_3D]]
// CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
@@ -18,7 +18,7 @@
// -----
module {
- // CHECK:spv.module "Logical" "GLSL450"
+ // CHECK:spv.module Logical GLSL450
// CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
// CHECK: spv.func [[FN:@broadcast_in_dim_scalar_3D]]
// CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/copy.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/copy.mlir
index eb7dead..f5f6bde 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/copy.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/copy.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -iree-index-computation -simplify-spirv-affine-exprs=false -convert-iree-to-spirv -verify-diagnostics -o - %s | IreeFileCheck %s
module {
- // CHECK:spv.module "Logical" "GLSL450"
+ // CHECK:spv.module Logical GLSL450
// CHECK-DAG: spv.globalVariable [[NUMWORKGROUPSVAR:@.*]] built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
// CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
// CHECK: spv.func [[FN:@simple_load_store]]
diff --git a/iree/samples/custom_modules/dialect/custom_ops.td b/iree/samples/custom_modules/dialect/custom_ops.td
index 05ff9c1..d721412 100644
--- a/iree/samples/custom_modules/dialect/custom_ops.td
+++ b/iree/samples/custom_modules/dialect/custom_ops.td
@@ -17,7 +17,7 @@
include "iree/compiler/Dialect/HAL/IR/HALBase.td"
include "iree/compiler/Dialect/IREE/IR/IREEBase.td"
-include "mlir/IR/OpBase.td"
+include "mlir/Interfaces/SideEffects.td"
def CUSTOM_Dialect : Dialect {
let name = "custom";
diff --git a/iree/vm/BUILD b/iree/vm/BUILD
index 4d9675f..c4e697b 100644
--- a/iree/vm/BUILD
+++ b/iree/vm/BUILD
@@ -104,6 +104,7 @@
"@llvm-project//mlir:OpBaseTdFiles",
"@llvm-project//mlir:include/mlir/Interfaces/CallInterfaces.td",
"@llvm-project//mlir:include/mlir/Interfaces/ControlFlowInterfaces.td",
+ "@llvm-project//mlir:include/mlir/Interfaces/SideEffects.td",
],
)
