Integrate llvm-project at 9816fada1667ecc3343d285295c95848826693af (#8612) * Reset third_party/llvm-project: 9816fada1667ecc3343d285295c95848826693af (2022-03-18 10:50:55 +0100): [bazel] Port a954ade8ed41 * Cherry-picked https://github.com/llvm/llvm-project/commit/3b74aac29c97b587c30d940e00b338af384ede59 additionally to fix bufferization issues * TensorFlow commit: c9ddfac348b1f423e1463ae78b046514a8b03a48 * Cherry-picked https://github.com/tensorflow/tensorflow/commit/9642bbdbb6b27443d339b714f233272eb74ca637 additionally to fix leftover `dump()` calls for debugging * MHLO commit: 467cd37703dc0c4195ce6351617ef320bb60e927 * Updated external model registration * Updated FuncOp usages * The attribute carrying type information now should be `function_type`, instead of `type`. * `mlir::FuncOp` is now in `mlir::func::FuncOp`. * `builtin.func` in assembly should be `func.func` now. * Updated MHLO attribute assembly for `comparison_direction`, `precision`, etc. * Fixed TFL/StripMetadata.cpp to avoid dropping `tf_saved_model` attributes for ABI information
diff --git a/integrations/tensorflow/WORKSPACE b/integrations/tensorflow/WORKSPACE index c947831..8284510 100644 --- a/integrations/tensorflow/WORKSPACE +++ b/integrations/tensorflow/WORKSPACE
@@ -7,7 +7,7 @@ load("@bazel_tools//tools/build_defs/repo:git.bzl", "git_repository") -TENSORFLOW_COMMIT = "05f17fca35623f4ab6d275ed95f0e1363c939f73" +TENSORFLOW_COMMIT = "9d9e78f6735a933b42febaafe2f303eec25f70ee" git_repository( name = "org_tensorflow",
diff --git a/integrations/tensorflow/iree-dialects/BUILD b/integrations/tensorflow/iree-dialects/BUILD index 3f47bd6..3280a69 100644 --- a/integrations/tensorflow/iree-dialects/BUILD +++ b/integrations/tensorflow/iree-dialects/BUILD
@@ -631,6 +631,7 @@ "@llvm-project//mlir:ArithmeticDialect", "@llvm-project//mlir:BufferizationDialect", "@llvm-project//mlir:BufferizationTransforms", + "@llvm-project//mlir:FuncDialect", "@llvm-project//mlir:LinalgOps", "@llvm-project//mlir:LLVMDialect", "@llvm-project//mlir:PDLDialect",
diff --git a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/PassDetail.h b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/PassDetail.h index 3f3fe9b..417a582 100644 --- a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/PassDetail.h +++ b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/PassDetail.h
@@ -7,6 +7,7 @@ #ifndef IREE_DIALECTS_DIALECT_LINALGEXT_TRANSFORMS_PASS_DETAIL_H_ #define IREE_DIALECTS_DIALECT_LINALGEXT_TRANSFORMS_PASS_DETAIL_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir {
diff --git a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.h b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.h index febec87..98e37ac 100644 --- a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.h +++ b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.h
@@ -7,6 +7,7 @@ #ifndef IREE_DIALECTS_DIALECT_LINALGEXT_TRANSFORMS_PASSES_H_ #define IREE_DIALECTS_DIALECT_LINALGEXT_TRANSFORMS_PASSES_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir { @@ -14,9 +15,9 @@ namespace IREE { namespace LinalgExt { -std::unique_ptr<OperationPass<FuncOp>> createTiledOpInterfaceTilingPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createTiledOpInterfaceTilingPass(); -std::unique_ptr<OperationPass<FuncOp>> createLinalgExtToLoopsPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createLinalgExtToLoopsPass(); std::unique_ptr<OperationPass<>> createPadContractionToBlockSizePass();
diff --git a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.td b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.td index 54a0484..95ae571 100644 --- a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.td +++ b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.td
@@ -10,13 +10,13 @@ include "mlir/Pass/PassBase.td" def LinalgExtToLoops : - Pass<"iree-linalg-ext-to-loops", "FuncOp"> { + Pass<"iree-linalg-ext-to-loops", "func::FuncOp"> { let summary = "Convert LinalgExt ops to loops and Linalg ops."; let constructor = "mlir::iree_compiler::IREE::LinalgExt::createLinalgExtToLoopsPass()"; } def TiledOpInterfaceTiling : - Pass<"iree-linalg-ext-tile", "FuncOp"> { + Pass<"iree-linalg-ext-tile", "func::FuncOp"> { let summary = "Test pass for tiling using TiledOpInterface"; let constructor = "mlir::iree_compiler::IREE::LinalgExt::createTiledOpInterfaceTilingPass()"; }
diff --git a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgTransform/LinalgTransformOps.h b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgTransform/LinalgTransformOps.h index 6eda492..7683ae0 100644 --- a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgTransform/LinalgTransformOps.h +++ b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/LinalgTransform/LinalgTransformOps.h
@@ -11,6 +11,7 @@ #include "iree-dialects/Dialect/LinalgExt/IR/LinalgExtOps.h" #include "iree-dialects/Dialect/LinalgTransform/TrackingListener.h" #include "iree-dialects/Dialect/LinalgTransform/TransformOpInterface.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/PDL/IR/PDLTypes.h" #include "mlir/IR/BuiltinAttributes.h"
diff --git a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.h b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.h index 905f3a3..c038c0b 100644 --- a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.h +++ b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.h
@@ -12,10 +12,12 @@ #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/Dialect.h" +#include "mlir/IR/FunctionInterfaces.h" #include "mlir/IR/OpDefinition.h" #include "mlir/IR/OpImplementation.h" #include "mlir/IR/PatternMatch.h" #include "mlir/IR/SymbolTable.h" +#include "mlir/Interfaces/CallInterfaces.h" #include "mlir/Interfaces/ControlFlowInterfaces.h" #include "mlir/Interfaces/SideEffectInterfaces.h"
diff --git a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.td b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.td index bc5b181..9b4aca3 100644 --- a/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.td +++ b/integrations/tensorflow/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.td
@@ -124,7 +124,7 @@ }]; let arguments = (ins SymbolNameAttr:$sym_name, - TypeAttr:$type, + TypeAttr:$function_type, OptionalAttr<StrArrayAttr>:$arg_names, OptionalAttr<StrArrayAttr>:$free_vars, OptionalAttr<StrArrayAttr>:$cell_vars, @@ -143,21 +143,21 @@ } /// Returns the type of this function. - FunctionType getType() { + FunctionType getFunctionType() { return getOperation()->getAttrOfType<TypeAttr>(getTypeAttrName()) .getValue() .cast<FunctionType>(); } /// Returns the argument types of this function. - ArrayRef<Type> getArgumentTypes() { return getType().getInputs(); } + ArrayRef<Type> getArgumentTypes() { return getFunctionType().getInputs(); } /// Returns the result types of this function. - ArrayRef<Type> getResultTypes() { return getType().getResults(); } + ArrayRef<Type> getResultTypes() { return getFunctionType().getResults(); } /// Returns the python return type of the function (second return type). Type getPyReturnType() { - return getType().getResult(1); + return getFunctionType().getResult(1); } /// Hook for Trait::FunctionLike, called after verifying that the 'type' @@ -167,7 +167,7 @@ Region *getCallableRegion() { return &body(); } ArrayRef<Type> getCallableResults() { - return getType().getResults(); + return getFunctionType().getResults(); } /// Defines SymbolOpInterface::isDeclaration(). @@ -177,8 +177,8 @@ }]; let builders = [ - OpBuilder<(ins "StringAttr":$name, "FunctionType":$type), [{ - build($_builder, $_state, name, TypeAttr::get(type), + OpBuilder<(ins "StringAttr":$name, "FunctionType":$function_type), [{ + build($_builder, $_state, name, TypeAttr::get(function_type), nullptr, nullptr, nullptr, nullptr); }]> ];
diff --git a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/IR/TiledOpInterface.cpp b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/IR/TiledOpInterface.cpp index e509489..604e498 100644 --- a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/IR/TiledOpInterface.cpp +++ b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/IR/TiledOpInterface.cpp
@@ -305,9 +305,12 @@ DialectRegistry ®istry) { LLVM_DEBUG( { llvm::dbgs() << "Adding external models of tiled op interface\n"; }); - registry - .addOpInterface<tensor::ExtractSliceOp, ExtractSliceTiledOpInterface>(); - registry.addOpInterface<tensor::InsertSliceOp, InsertSliceTiledOpInterface>(); + + registry.addExtension(+[](MLIRContext *ctx, tensor::TensorDialect *dialect) { + tensor::ExtractSliceOp::attachInterface<ExtractSliceTiledOpInterface>(*ctx); + tensor::InsertSliceOp::attachInterface<InsertSliceTiledOpInterface>(*ctx); + }); + // TODO(ravishankarm): Needs custom PadTiledOpInterface or equiv. // registry.addOpInterface<tensor::PadOp, // ForwardToTilingInterface<tensor::PadOp>>();
diff --git a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/LinalgExtBufferization.cpp b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/LinalgExtBufferization.cpp index 2b8f8ec..9b3c2fb 100644 --- a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/LinalgExtBufferization.cpp +++ b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/LinalgExtBufferization.cpp
@@ -6,11 +6,11 @@ #include "iree-dialects/Dialect/LinalgExt/LinalgExtBufferization.h" -#include <mlir/IR/BuiltinOps.h> - +#include "iree-dialects/Dialect/LinalgExt/IR/LinalgExtDialect.h" #include "iree-dialects/Dialect/LinalgExt/IR/LinalgExtOps.h" #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" +#include "mlir/IR/BuiltinOps.h" #include "mlir/IR/PatternMatch.h" using namespace mlir; @@ -22,7 +22,6 @@ using bufferization::BufferRelation; using bufferization::getMemRefType; using bufferization::replaceOpWithBufferizedValues; -using bufferization::replaceOpWithNewBufferizedOp; using tensor::ExtractSliceOp; /// Return the destinations that an InParallelOp is inserting into. One per @@ -341,9 +340,12 @@ void mlir::iree_compiler::IREE::LinalgExt:: registerBufferizableOpInterfaceExternalModels(DialectRegistry ®istry) { - registry.addOpInterface<InParallelOp, InParallelOpInterface>(); - registry - .addOpInterface<PerformConcurrentlyOp, PerformConcurrentlyOpInterface>(); - registry - .addOpInterface<ParallelInsertSliceOp, ParallelInsertSliceOpInterface>(); + registry.addExtension( + +[](MLIRContext *ctx, LinalgExt::IREELinalgExtDialect *dialect) { + InParallelOp::attachInterface<InParallelOpInterface>(*ctx); + PerformConcurrentlyOp::attachInterface<PerformConcurrentlyOpInterface>( + *ctx); + ParallelInsertSliceOp::attachInterface<ParallelInsertSliceOpInterface>( + *ctx); + }); }
diff --git a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/TilingExternalModels.cpp b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/TilingExternalModels.cpp index 174d4ff..0758d7f 100644 --- a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/TilingExternalModels.cpp +++ b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/TilingExternalModels.cpp
@@ -157,23 +157,25 @@ } // namespace template <typename OpType> -void registerOne(DialectRegistry ®istry) { - registry.addOpInterface<OpType, LinalgOpTilingInterface<OpType>>(); +void registerOne(MLIRContext *ctx) { + OpType::template attachInterface<LinalgOpTilingInterface<OpType>>(*ctx); } /// Variadic helper function. template <typename... OpTypes> -void registerAll(DialectRegistry ®istry) { +void registerAll(MLIRContext *ctx) { // FIXME: In c++17 this can be simplified by using 'fold expressions'. - (void)std::initializer_list<int>{0, (registerOne<OpTypes>(registry), 0)...}; + (void)std::initializer_list<int>{0, (registerOne<OpTypes>(ctx), 0)...}; } #define GET_OP_LIST void mlir::iree_compiler::IREE::LinalgExt:: registerTilingInterfaceExternalModels(DialectRegistry ®istry) { - registerOne<linalg::GenericOp>(registry); - registerAll< + registry.addExtension(+[](MLIRContext *ctx, linalg::LinalgDialect *dialect) { + registerOne<linalg::GenericOp>(ctx); + registerAll< #include "mlir/Dialect/Linalg/IR/LinalgStructuredOps.cpp.inc" - >(registry); + >(ctx); + }); }
diff --git a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgTransform/IR/LinalgTransformOps.cpp b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgTransform/IR/LinalgTransformOps.cpp index e28dcd3..03d39ac 100644 --- a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgTransform/IR/LinalgTransformOps.cpp +++ b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgTransform/IR/LinalgTransformOps.cpp
@@ -31,6 +31,7 @@ #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" #include "mlir/Dialect/Bufferization/Transforms/OneShotAnalysis.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/Linalg/ComprehensiveBufferize/ModuleBufferization.h" #include "mlir/Dialect/Linalg/Passes.h" @@ -581,7 +582,7 @@ // Perform buffer-level hoistings. state.getTopLevel()->walk( - [&](FuncOp funcOp) { hoistRedundantVectorTransfers(funcOp); }); + [&](func::FuncOp funcOp) { hoistRedundantVectorTransfers(funcOp); }); return success(); } @@ -597,8 +598,8 @@ // the end. Keep module-level for now. PassManager pm(getContext()); - pm.addNestedPass<FuncOp>(createConvertVectorToSCFPass()); - pm.addNestedPass<FuncOp>(createConvertLinalgToLoopsPass()); + pm.addNestedPass<func::FuncOp>(createConvertVectorToSCFPass()); + pm.addNestedPass<func::FuncOp>(createConvertLinalgToLoopsPass()); if (enable_async()) { pm.addPass(createAsyncToAsyncRuntimePass()); pm.addPass(createAsyncRuntimeRefCountingPass()); @@ -618,7 +619,7 @@ .enableAMX(enable_amx()) .enableX86Vector(enable_x86vector()))); // clang-format on - pm.addNestedPass<FuncOp>(createConvertMathToLLVMPass()); + pm.addNestedPass<func::FuncOp>(createConvertMathToLLVMPass()); pm.addPass(createMemRefToLLVMPass()); if (enable_async()) pm.addPass(createConvertAsyncToLLVMPass()); @@ -631,7 +632,9 @@ // FIXME: this is a terrible hack! state.getTopLevel()->walk([](LLVM::LLVMFuncOp funcOp) { for (int64_t i = 0; i < funcOp.getNumArguments(); ++i) { - if (!funcOp.getType().getParamType(i).isa<LLVM::LLVMPointerType>()) + if (!funcOp.getFunctionType() + .getParamType(i) + .isa<LLVM::LLVMPointerType>()) continue; funcOp.setArgAttr(i, "llvm.noalias", UnitAttr::get(funcOp.getContext())); } @@ -760,15 +763,15 @@ return executeRegionOp; } -static FailureOr<FuncOp> outlineLoop(scf::ForOp loop, StringRef funcName, - transform::TransformState &state) { +static FailureOr<func::FuncOp> outlineLoop(scf::ForOp loop, StringRef funcName, + transform::TransformState &state) { PatternRewriterListener rewriter(loop->getContext()); auto &listener = state.getExtension<TrackingListener>(); rewriter.addListener(&listener); Location loc = loop.getLoc(); scf::ExecuteRegionOp exec = outlineInExecuteRegion(rewriter, loop); assert(exec && "failed to produce execute_region"); - FailureOr<FuncOp> outlined = + FailureOr<func::FuncOp> outlined = outlineSingleBlockRegion(rewriter, loc, exec.getRegion(), funcName); if (failed(listener.checkErrorState())) return failure(); @@ -781,7 +784,7 @@ SmallVector<Operation *> resultVector; auto res = applyTransformToEach(state.getPayloadOps(target()), resultVector, - [&](scf::ForOp loop) -> FailureOr<FuncOp> { + [&](scf::ForOp loop) -> FailureOr<func::FuncOp> { return outlineLoop(loop, func_name(), state); }); if (failed(res))
diff --git a/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/IR/PyDMOps.cpp b/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/IR/PyDMOps.cpp index b281874..f4ad03a 100644 --- a/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/IR/PyDMOps.cpp +++ b/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/IR/PyDMOps.cpp
@@ -567,9 +567,7 @@ } void PyFuncOp::print(OpAsmPrinter &p) { - FunctionType fnType = getType(); - function_interface_impl::printFunctionOp( - p, *this, fnType.getInputs(), /*isVariadic=*/false, fnType.getResults()); + function_interface_impl::printFunctionOp(p, *this, /*isVariadic=*/false); } //===----------------------------------------------------------------------===// @@ -764,7 +762,7 @@ << "' does not reference a valid function"; // Verify that the operand and result types match the callee. - auto fnType = fn.getType(); + auto fnType = fn.getFunctionType(); if (fnType.getNumInputs() != getNumOperands()) return emitOpError("incorrect number of operands for callee");
diff --git a/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/Transforms/Optimize/FixateWeakNumeric.cpp b/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/Transforms/Optimize/FixateWeakNumeric.cpp index b6e5b51..5d4be33 100644 --- a/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/Transforms/Optimize/FixateWeakNumeric.cpp +++ b/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/Transforms/Optimize/FixateWeakNumeric.cpp
@@ -41,7 +41,7 @@ // Special cases for operations. if (auto funcOp = llvm::dyn_cast<PYDM::FuncOp>(op)) { - FunctionType existingFt = funcOp.getType(); + FunctionType existingFt = funcOp.getFunctionType(); FunctionType newFt = convertFunctionType(existingFt); if (newFt != existingFt) { funcOp.setType(newFt);
diff --git a/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/Transforms/ToIREE/LoweringPatterns.cpp b/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/Transforms/ToIREE/LoweringPatterns.cpp index 9d92711..2e14e37 100644 --- a/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/Transforms/ToIREE/LoweringPatterns.cpp +++ b/integrations/tensorflow/iree-dialects/lib/Dialect/PyDM/Transforms/ToIREE/LoweringPatterns.cpp
@@ -637,7 +637,7 @@ LogicalResult matchAndRewrite(PYDM::FuncOp srcOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - FunctionType srcFuncType = srcOp.getType(); + FunctionType srcFuncType = srcOp.getFunctionType(); TypeConverter::SignatureConversion signatureConversion( srcOp.getNumArguments()); @@ -839,7 +839,7 @@ auto parentFunc = srcOp->getParentOfType<mlir::FuncOp>(); if (!parentFunc) return rewriter.notifyMatchFailure(srcOp, "not contained by a func"); - Type convertedReturnType = parentFunc.getType().getResult(1); + Type convertedReturnType = parentFunc.getFunctionType().getResult(1); // Split the entry block. Block *entryBlock = rewriter.getInsertionBlock();
diff --git a/integrations/tensorflow/iree-dialects/python/iree/compiler/dialects/_iree_pydm_ops_ext.py b/integrations/tensorflow/iree-dialects/python/iree/compiler/dialects/_iree_pydm_ops_ext.py index 5155972..7fc8653 100644 --- a/integrations/tensorflow/iree-dialects/python/iree/compiler/dialects/_iree_pydm_ops_ext.py +++ b/integrations/tensorflow/iree-dialects/python/iree/compiler/dialects/_iree_pydm_ops_ext.py
@@ -18,7 +18,7 @@ @property def type(self): - return ir.FunctionType(ir.TypeAttr(self.attributes["type"]).value) + return ir.FunctionType(ir.TypeAttr(self.attributes["function_type"]).value) @property def py_return_type(self) -> ir.Type:
diff --git a/integrations/tensorflow/iree-dialects/python/iree/compiler/dialects/iree_pydm/importer/importer.py b/integrations/tensorflow/iree-dialects/python/iree/compiler/dialects/iree_pydm/importer/importer.py index 68b1227..fae6591 100644 --- a/integrations/tensorflow/iree-dialects/python/iree/compiler/dialects/iree_pydm/importer/importer.py +++ b/integrations/tensorflow/iree-dialects/python/iree/compiler/dialects/iree_pydm/importer/importer.py
@@ -85,7 +85,7 @@ context=ic.context) f_op = d.FuncOp( ir.StringAttr.get(symbol), - type=ir.TypeAttr.get(ir_f_type), + function_type=ir.TypeAttr.get(ir_f_type), arg_names=f_arg_names, free_vars=f_var_names, cell_vars=ir.ArrayAttr.get([]),
diff --git a/integrations/tensorflow/iree-dialects/test/Transforms/test-listener-cse.mlir b/integrations/tensorflow/iree-dialects/test/Transforms/test-listener-cse.mlir index 13d2994..ebfc348 100644 --- a/integrations/tensorflow/iree-dialects/test/Transforms/test-listener-cse.mlir +++ b/integrations/tensorflow/iree-dialects/test/Transforms/test-listener-cse.mlir
@@ -233,7 +233,7 @@ %0 = arith.constant 1 : i32 // CHECK-NEXT: @nested_func - builtin.func @nested_func() { + func.func @nested_func() { // CHECK-NEXT: arith.constant 1 %foo = arith.constant 1 : i32 "foo.yield"(%foo) : (i32) -> ()
diff --git a/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.cpp b/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.cpp index d2cd879..d26e8fc 100644 --- a/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.cpp +++ b/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.cpp
@@ -25,7 +25,7 @@ // Import pipelines should end with canonicalization because they may have // access to dialects and patterns that the core compiler does not. - pm.addNestedPass<FuncOp>(mlir::createCanonicalizerPass()); + pm.addNestedPass<func::FuncOp>(mlir::createCanonicalizerPass()); } void registerMHLOImportPassPipeline() {
diff --git a/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.h b/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.h index 670104a..8930e85 100644 --- a/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.h +++ b/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.h
@@ -7,6 +7,7 @@ #ifndef IREE_INTEGRATIONS_TENSORFLOW_IREE_TF_COMPILER_MHLO_PASSES_H_ #define IREE_INTEGRATIONS_TENSORFLOW_IREE_TF_COMPILER_MHLO_PASSES_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir { @@ -27,7 +28,7 @@ // Annotates an appropriate iree.abi attribute on public functions that // operate exclusively on tensor types. This corresponds to the expectations // of MHLO and is suitable for such programs. -std::unique_ptr<OperationPass<FuncOp>> createEmitDefaultIREEABIPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createEmitDefaultIREEABIPass(); //===----------------------------------------------------------------------===// // Registration
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/Passes.h b/integrations/tensorflow/iree_tf_compiler/TF/Passes.h index 6a3712c..9e9d6c3 100644 --- a/integrations/tensorflow/iree_tf_compiler/TF/Passes.h +++ b/integrations/tensorflow/iree_tf_compiler/TF/Passes.h
@@ -7,6 +7,7 @@ #ifndef IREE_INTEGRATIONS_TENSORFLOW_IREE_TF_COMPILER_TF_PASSES_H_ #define IREE_INTEGRATIONS_TENSORFLOW_IREE_TF_COMPILER_TF_PASSES_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir { @@ -48,14 +49,14 @@ std::unique_ptr<OperationPass<ModuleOp>> createPropagateResourceCastsPass(); // Strips tf.Assert ops. -std::unique_ptr<OperationPass<FuncOp>> createStripAssertsPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createStripAssertsPass(); // Strips all TF-related attributes; none are needed by IREE. std::unique_ptr<OperationPass<ModuleOp>> createStripModuleMetadataPass(); -std::unique_ptr<OperationPass<FuncOp>> createStripFunctionMetadataPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createStripFunctionMetadataPass(); // Validates whether any Tensorflow operations remain. -std::unique_ptr<OperationPass<FuncOp>> createVerifyFullyConvertedPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createVerifyFullyConvertedPass(); //===----------------------------------------------------------------------===// // Patterns
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/test/convert_to_mhlo.mlir b/integrations/tensorflow/iree_tf_compiler/TF/test/convert_to_mhlo.mlir index dbc1ec7..87213e8 100644 --- a/integrations/tensorflow/iree_tf_compiler/TF/test/convert_to_mhlo.mlir +++ b/integrations/tensorflow/iree_tf_compiler/TF/test/convert_to_mhlo.mlir
@@ -2,7 +2,10 @@ // CHECK-LABEL: @sigmoid func @sigmoid(%arg0: tensor<2xf32>) -> tensor<2xf32> { - // CHECK-DAG: [[HALF:%.+]] = mhlo.constant dense<5.000000e-01> : tensor<2xf32> + // CHECK-DAG: [[SCALAR:%.+]] = mhlo.constant dense<5.000000e-01> : tensor<f32> + // CHECK-DAG: [[SHAPE_OF:%.+]] = shape.shape_of %arg0 : tensor<2xf32> -> tensor<1xindex> + // CHECK-DAG: [[SHAPE_VAL:%.+]] = shape.to_extent_tensor [[SHAPE_OF]] : tensor<1xindex> -> tensor<1xindex> + // CHECK-DAG: [[HALF:%.+]] = "mhlo.dynamic_broadcast_in_dim"([[SCALAR]], [[SHAPE_VAL]]) {broadcast_dimensions = dense<> : tensor<0xi64>} : (tensor<f32>, tensor<1xindex>) -> tensor<2xf32> // CHECK-DAG: [[R1:%.+]] = mhlo.multiply %arg0, [[HALF]] : tensor<2xf32> // CHECK-DAG: [[R2:%.+]] = "mhlo.tanh"([[R1]]) : (tensor<2xf32>) -> tensor<2xf32> // CHECK-DAG: [[R3:%.+]] = mhlo.multiply [[R2]], [[HALF]] : tensor<2xf32> @@ -13,7 +16,7 @@ // CHECK-LABEL: @sigmoid_complex func @sigmoid_complex(%arg0: tensor<2xcomplex<f32>>) -> tensor<2xcomplex<f32>> { - // CHECK: [[R0:%.+]] = mhlo.constant dense<(5.000000e-01,0.000000e+00)> : tensor<2xcomplex<f32>> + // CHECK: [[R0:%.+]] = mhlo.constant dense<(5.000000e-01,0.000000e+00)> : tensor<complex<f32>> // CHECK-NOT: tf.Sigmoid %0 = "tf.Sigmoid"(%arg0) : (tensor<2xcomplex<f32>>) -> tensor<2xcomplex<f32>> return %0 : tensor<2xcomplex<f32>> @@ -22,7 +25,8 @@ // CHECK-LABEL: @sigmoid_unranked func @sigmoid_unranked(%arg0: tensor<*xf32>) -> tensor<*xf32> { // CHECK-DAG: [[SCALAR:%.+]] = mhlo.constant dense<5.000000e-01> : tensor<f32> - // CHECK-DAG: [[SHAPE_VAL:%.+]] = shape.shape_of %arg0 : tensor<*xf32> -> tensor<?xindex> + // CHECK-DAG: [[SHAPE_OF:%.+]] = shape.shape_of %arg0 : tensor<*xf32> -> tensor<?xindex> + // CHECK-DAG: [[SHAPE_VAL:%.+]] = shape.to_extent_tensor [[SHAPE_OF]] : tensor<?xindex> -> tensor<?xindex> // CHECK-DAG: [[HALF:%.+]] = "mhlo.dynamic_broadcast_in_dim"([[SCALAR]], [[SHAPE_VAL]]) {broadcast_dimensions = dense<> : tensor<0xi64>} : (tensor<f32>, tensor<?xindex>) -> tensor<*xf32> // CHECK-DAG: [[R1:%.+]] = mhlo.multiply %arg0, [[HALF]] : tensor<*xf32> // CHECK-DAG: [[R2:%.+]] = "mhlo.tanh"([[R1]]) : (tensor<*xf32>) -> tensor<*xf32>
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/test/saved_model_to_iree_abi.mlir b/integrations/tensorflow/iree_tf_compiler/TF/test/saved_model_to_iree_abi.mlir index 029e580..bf5a2e7 100644 --- a/integrations/tensorflow/iree_tf_compiler/TF/test/saved_model_to_iree_abi.mlir +++ b/integrations/tensorflow/iree_tf_compiler/TF/test/saved_model_to_iree_abi.mlir
@@ -13,7 +13,7 @@ // CHECK: func private @__inference_binary_func_70 // CHECK-NOT: tf_saved_model builtin.module @binary_func attributes {tf.versions = {bad_consumers = [], min_consumer = 12 : i32, producer = 729 : i32}, tf_saved_model.semantics} { - builtin.func @__inference_binary_func_70(%arg0: tensor<16xf32> {tf._user_specified_name = "a", tf_saved_model.index_path = [0]}, %arg1: tensor<16xf32> {tf._user_specified_name = "b", tf_saved_model.index_path = [1]}) -> (tensor<16xf32> {tf_saved_model.index_path = [0]}, tensor<16xf32> {tf_saved_model.index_path = [1]}) attributes {tf._construction_context = "kEagerRuntime", tf._input_shapes = [#tf_type.shape<16>, #tf_type.shape<16>], tf_saved_model.exported_names = ["binary_func"]} { + func.func @__inference_binary_func_70(%arg0: tensor<16xf32> {tf._user_specified_name = "a", tf_saved_model.index_path = [0]}, %arg1: tensor<16xf32> {tf._user_specified_name = "b", tf_saved_model.index_path = [1]}) -> (tensor<16xf32> {tf_saved_model.index_path = [0]}, tensor<16xf32> {tf_saved_model.index_path = [1]}) attributes {tf._construction_context = "kEagerRuntime", tf._input_shapes = [#tf_type.shape<16>, #tf_type.shape<16>], tf_saved_model.exported_names = ["binary_func"]} { %0 = "tf.Identity"(%arg0) {device = ""} : (tensor<16xf32>) -> tensor<16xf32> %1 = "tf.Identity"(%arg1) {device = ""} : (tensor<16xf32>) -> tensor<16xf32> return %0, %1 : tensor<16xf32>, tensor<16xf32> @@ -31,7 +31,7 @@ // CHECK: func private @__inference_unary_func_240 // CHECK-NOT: tf_saved_model builtin.module @unary_func attributes {tf.versions = {bad_consumers = [], min_consumer = 12 : i32, producer = 729 : i32}, tf_saved_model.semantics} { - builtin.func @__inference_unary_func_240(%arg0: tensor<16xf32> {tf._user_specified_name = "a", tf_saved_model.index_path = [0]}) -> (tensor<16xf32> {tf_saved_model.index_path = []}) attributes {tf._construction_context = "kEagerRuntime", tf._input_shapes = [#tf_type.shape<16>], tf_saved_model.exported_names = ["unary_func"]} { + func.func @__inference_unary_func_240(%arg0: tensor<16xf32> {tf._user_specified_name = "a", tf_saved_model.index_path = [0]}) -> (tensor<16xf32> {tf_saved_model.index_path = []}) attributes {tf._construction_context = "kEagerRuntime", tf._input_shapes = [#tf_type.shape<16>], tf_saved_model.exported_names = ["unary_func"]} { %0 = "tf.Identity"(%arg0) {device = ""} : (tensor<16xf32>) -> tensor<16xf32> return %0 : tensor<16xf32> } @@ -50,7 +50,7 @@ // CHECK: func private @__inference_return_list_260 // CHECK-NOT: tf_saved_model builtin.module @return_list attributes {tf.versions = {bad_consumers = [], min_consumer = 12 : i32, producer = 729 : i32}, tf_saved_model.semantics} { - builtin.func @__inference_return_list_260(%arg0: tensor<16xf32> {tf._user_specified_name = "a", tf_saved_model.index_path = [0]}, %arg1: tensor<16xf32> {tf._user_specified_name = "b", tf_saved_model.index_path = [1]}) -> (tensor<16xf32> {tf_saved_model.index_path = [0]}, tensor<16xf32> {tf_saved_model.index_path = [1]}) attributes {tf._construction_context = "kEagerRuntime", tf._input_shapes = [#tf_type.shape<16>, #tf_type.shape<16>], tf_saved_model.exported_names = ["return_list"]} { + func.func @__inference_return_list_260(%arg0: tensor<16xf32> {tf._user_specified_name = "a", tf_saved_model.index_path = [0]}, %arg1: tensor<16xf32> {tf._user_specified_name = "b", tf_saved_model.index_path = [1]}) -> (tensor<16xf32> {tf_saved_model.index_path = [0]}, tensor<16xf32> {tf_saved_model.index_path = [1]}) attributes {tf._construction_context = "kEagerRuntime", tf._input_shapes = [#tf_type.shape<16>, #tf_type.shape<16>], tf_saved_model.exported_names = ["return_list"]} { %0 = "tf.Identity"(%arg0) {device = ""} : (tensor<16xf32>) -> tensor<16xf32> %1 = "tf.Identity"(%arg1) {device = ""} : (tensor<16xf32>) -> tensor<16xf32> return %0, %1 : tensor<16xf32>, tensor<16xf32> @@ -101,7 +101,7 @@ // CHECK: func private @__inference_dict_nest_190 // CHECK-NOT: tf_saved_model builtin.module @dict_nest attributes {tf.versions = {bad_consumers = [], min_consumer = 12 : i32, producer = 729 : i32}, tf_saved_model.semantics} { - builtin.func @__inference_dict_nest_190( + func.func @__inference_dict_nest_190( %arg0: tensor<16xf32> {tf._user_specified_name = "mapping", tf_saved_model.index_path = [0, "dict", "a"]}, %arg1: tensor<16xf32> {tf._user_specified_name = "mapping", tf_saved_model.index_path = [0, "dict", "b"]}, %arg2: tensor<16xf32> {tf._user_specified_name = "mapping", tf_saved_model.index_path = [0, "list", 0]}, @@ -120,7 +120,7 @@ // CHECK: func @dict_nest(%arg0: !iree_input.buffer_view, %arg1: !iree_input.buffer_view, %arg2: !iree_input.buffer_view) -> !iree_input.list<!iree_input.variant> // CHECK-SAME{LITERAL}: iree.abi = "{\22a\22:[[\22named\22,\22a\22,[\22ndarray\22,\22f32\22,1,16]],[\22named\22,\22b\22,[\22ndarray\22,\22f32\22,1,16]],[\22named\22,\22scalar\22,[\22ndarray\22,\22f32\22,0]]],\22r\22:[[\22sdict\22,[\22dict\22,[\22sdict\22,[\22a\22,[\22ndarray\22,\22f32\22,1,16]],[\22b\22,[\22ndarray\22,\22f32\22,1,16]],[\22scalar\22,[\22ndarray\22,\22f32\22,0]]]]]],\22v\22:1}" builtin.module @kwargs attributes {tf.versions = {bad_consumers = [], min_consumer = 12 : i32, producer = 729 : i32}, tf_saved_model.semantics} { - builtin.func @__inference_dict_nest_190( + func.func @__inference_dict_nest_190( %arg0: tensor<16xf32> {tf_saved_model.index_path = ["a"]}, %arg1: tensor<16xf32> {tf_saved_model.index_path = ["b"]}, %arg2: tensor<f32> {tf._user_specified_name = "scalar", tf_saved_model.index_path = ["scalar"]}) ->
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/test/strip_asserts.mlir b/integrations/tensorflow/iree_tf_compiler/TF/test/strip_asserts.mlir index 2711fd4..d784b98 100644 --- a/integrations/tensorflow/iree_tf_compiler/TF/test/strip_asserts.mlir +++ b/integrations/tensorflow/iree_tf_compiler/TF/test/strip_asserts.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-tf-opt -split-input-file -verify-diagnostics -pass-pipeline='builtin.func(iree-tf-strip-asserts)' %s | FileCheck %s +// RUN: iree-tf-opt -split-input-file -verify-diagnostics -pass-pipeline='func.func(iree-tf-strip-asserts)' %s | FileCheck %s // CHECK-LABEL: @asserts // CHECK-NOT: tf.Assert
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/test/strip_metadata.mlir b/integrations/tensorflow/iree_tf_compiler/TF/test/strip_metadata.mlir index db23522..1cc7e83 100644 --- a/integrations/tensorflow/iree_tf_compiler/TF/test/strip_metadata.mlir +++ b/integrations/tensorflow/iree_tf_compiler/TF/test/strip_metadata.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-tf-opt -split-input-file -verify-diagnostics -pass-pipeline='iree-tf-strip-module-metadata,builtin.func(iree-tf-strip-function-metadata)' %s | FileCheck %s +// RUN: iree-tf-opt -split-input-file -verify-diagnostics -pass-pipeline='iree-tf-strip-module-metadata,func.func(iree-tf-strip-function-metadata)' %s | FileCheck %s // CHECK-LABEL: @tf_module // CHECK-NOT: attributes
diff --git a/integrations/tensorflow/iree_tf_compiler/TFL/LowerGlobalTensors.cpp b/integrations/tensorflow/iree_tf_compiler/TFL/LowerGlobalTensors.cpp index f7765a1..ecc9ab6 100644 --- a/integrations/tensorflow/iree_tf_compiler/TFL/LowerGlobalTensors.cpp +++ b/integrations/tensorflow/iree_tf_compiler/TFL/LowerGlobalTensors.cpp
@@ -37,7 +37,7 @@ DenseMap<StringRef, FuncOp> symNameToFunction; for (auto func : moduleOp.getOps<FuncOp>()) { - symNameToFunction[func.sym_name()] = func; + symNameToFunction[func.getSymName()] = func; } DenseMap<StringRef, DenseElementsAttr> sharedNameToConstant;
diff --git a/integrations/tensorflow/iree_tf_compiler/TFL/PassDetail.h b/integrations/tensorflow/iree_tf_compiler/TFL/PassDetail.h index 0d943d7..5cc38bc 100644 --- a/integrations/tensorflow/iree_tf_compiler/TFL/PassDetail.h +++ b/integrations/tensorflow/iree_tf_compiler/TFL/PassDetail.h
@@ -7,6 +7,7 @@ #ifndef IREE_INTEGRATIONS_TENSORFLOW_IREE_TF_COMPILER_TFL_PASS_DETAIL_H_ #define IREE_INTEGRATIONS_TENSORFLOW_IREE_TF_COMPILER_TFL_PASS_DETAIL_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir {
diff --git a/integrations/tensorflow/iree_tf_compiler/TFL/Passes.h b/integrations/tensorflow/iree_tf_compiler/TFL/Passes.h index bd836f6..f04f64d 100644 --- a/integrations/tensorflow/iree_tf_compiler/TFL/Passes.h +++ b/integrations/tensorflow/iree_tf_compiler/TFL/Passes.h
@@ -7,6 +7,7 @@ #ifndef IREE_INTEGRATIONS_TENSORFLOW_IREE_TF_COMPILER_TFL_PASSES_H_ #define IREE_INTEGRATIONS_TENSORFLOW_IREE_TF_COMPILER_TFL_PASSES_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir { @@ -30,17 +31,18 @@ // Converts TFLite attributes that are useful to corresponding IREE attributes. std::unique_ptr<OperationPass<ModuleOp>> createConvertModuleMetadataPass(); -std::unique_ptr<OperationPass<FuncOp>> createConvertFunctionMetadataPass(); +std::unique_ptr<OperationPass<func::FuncOp>> +createConvertFunctionMetadataPass(); // Lowers TFLite's global tensor operations to the Util dialect. std::unique_ptr<OperationPass<ModuleOp>> createLowerGlobalTensorsPass(); // Strips all leftover TFLite-related attributes; none are needed by IREE. std::unique_ptr<OperationPass<ModuleOp>> createStripModuleMetadataPass(); -std::unique_ptr<OperationPass<FuncOp>> createStripFunctionMetadataPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createStripFunctionMetadataPass(); // Validates whether any TFLite operations remain. -std::unique_ptr<OperationPass<FuncOp>> createVerifyFullyConvertedPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createVerifyFullyConvertedPass(); //===----------------------------------------------------------------------===// // Registration
diff --git a/integrations/tensorflow/iree_tf_compiler/TFL/Passes.td b/integrations/tensorflow/iree_tf_compiler/TFL/Passes.td index 43ab16a..70eaab7 100644 --- a/integrations/tensorflow/iree_tf_compiler/TFL/Passes.td +++ b/integrations/tensorflow/iree_tf_compiler/TFL/Passes.td
@@ -10,7 +10,7 @@ include "mlir/Pass/PassBase.td" def ConvertFunctionMetadata : - Pass<"iree-tflite-convert-function-metadata", "mlir::FuncOp"> { + Pass<"iree-tflite-convert-function-metadata", "mlir::func::FuncOp"> { let summary = "Converts TFLite attributes to IREE attributes on functions."; let constructor = "mlir::iree_integrations::TFL::createConvertFunctionMetadataPass()"; } @@ -34,7 +34,7 @@ } def StripFunctionMetadata : - Pass<"iree-tflite-strip-function-metadata", "mlir::FuncOp"> { + Pass<"iree-tflite-strip-function-metadata", "mlir::func::FuncOp"> { let summary = "Guarantees that functions used by tfl.call_once are retained."; let constructor = "mlir::iree_integrations::TFL::createStripFunctionMetadataPass()"; } @@ -46,7 +46,7 @@ } def VerifyFullyConverted : - Pass<"iree-tflite-verify-fully-converted", "mlir::FuncOp"> { + Pass<"iree-tflite-verify-fully-converted", "mlir::func::FuncOp"> { let summary = "Verifies that all TFLite frontend ops were converted and none remain."; let constructor = "mlir::iree_integrations::TFL::createVerifyFullyConvertedPass()"; }
diff --git a/integrations/tensorflow/iree_tf_compiler/TFL/RetainCallOnceFuncs.cpp b/integrations/tensorflow/iree_tf_compiler/TFL/RetainCallOnceFuncs.cpp index f6b9d94..c9a605e 100644 --- a/integrations/tensorflow/iree_tf_compiler/TFL/RetainCallOnceFuncs.cpp +++ b/integrations/tensorflow/iree_tf_compiler/TFL/RetainCallOnceFuncs.cpp
@@ -29,11 +29,11 @@ auto moduleOp = getOperation(); llvm::DenseMap<StringRef, FuncOp> funcMap; - for (auto func : moduleOp.getOps<mlir::FuncOp>()) { - funcMap[func.sym_name()] = func; + for (auto func : moduleOp.getOps<mlir::func::FuncOp>()) { + funcMap[func.getSymName()] = func; } - for (auto func : moduleOp.getOps<mlir::FuncOp>()) { + for (auto func : moduleOp.getOps<mlir::func::FuncOp>()) { for (auto callOnce : func.getOps<mlir::TFL::CallOnceOp>()) { auto callFunc = funcMap[callOnce.session_init_function()]; callOnce->setAttr("session_init_function_symbol",
diff --git a/integrations/tensorflow/iree_tf_compiler/TFL/StripMetadata.cpp b/integrations/tensorflow/iree_tf_compiler/TFL/StripMetadata.cpp index af1dbe6..181962c 100644 --- a/integrations/tensorflow/iree_tf_compiler/TFL/StripMetadata.cpp +++ b/integrations/tensorflow/iree_tf_compiler/TFL/StripMetadata.cpp
@@ -17,6 +17,9 @@ static bool isTFLAttr(NamedAttribute &namedAttr) { // NOTE: tflite mixes tf and tfl, for some reason. auto name = namedAttr.getName().strref(); + // Don't trim attributes from tf_saved_model---they carry ABI information. + if (name.startswith("tf_saved_model.")) return false; + if (name.startswith("tf.") || name.startswith("tf_") || name.startswith("tfl.") || name.startswith("tfl_")) { return true;
diff --git a/integrations/tensorflow/iree_tf_compiler/TFL/test/convert_metadata.mlir b/integrations/tensorflow/iree_tf_compiler/TFL/test/convert_metadata.mlir index 4a8825e..84752cd 100644 --- a/integrations/tensorflow/iree_tf_compiler/TFL/test/convert_metadata.mlir +++ b/integrations/tensorflow/iree_tf_compiler/TFL/test/convert_metadata.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt-tflite -split-input-file -pass-pipeline='iree-tflite-convert-module-metadata,builtin.func(iree-tflite-convert-function-metadata)' %s | FileCheck %s +// RUN: iree-opt-tflite -split-input-file -pass-pipeline='iree-tflite-convert-module-metadata,func.func(iree-tflite-convert-function-metadata)' %s | FileCheck %s module attributes {tfl.schema_version = 3 : i32} { // CHECK: func @main(
diff --git a/integrations/tensorflow/iree_tf_compiler/TFL/test/strip_metadata.mlir b/integrations/tensorflow/iree_tf_compiler/TFL/test/strip_metadata.mlir index 8cce458..b0b192c 100644 --- a/integrations/tensorflow/iree_tf_compiler/TFL/test/strip_metadata.mlir +++ b/integrations/tensorflow/iree_tf_compiler/TFL/test/strip_metadata.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt-tflite -split-input-file -verify-diagnostics -pass-pipeline='iree-tflite-strip-module-metadata,builtin.func(iree-tflite-strip-function-metadata)' %s | FileCheck %s +// RUN: iree-opt-tflite -split-input-file -verify-diagnostics -pass-pipeline='iree-tflite-strip-module-metadata,func.func(iree-tflite-strip-function-metadata)' %s | FileCheck %s // CHECK-LABEL: module { // CHECK-NOT: tf.schema_version
diff --git a/iree/compiler/Bindings/Native/Transforms/Passes.cpp b/iree/compiler/Bindings/Native/Transforms/Passes.cpp index 4ec5152..65dd8e4 100644 --- a/iree/compiler/Bindings/Native/Transforms/Passes.cpp +++ b/iree/compiler/Bindings/Native/Transforms/Passes.cpp
@@ -8,6 +8,7 @@ #include <memory> +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/PassOptions.h" #include "mlir/Pass/PassRegistry.h" #include "mlir/Transforms/Passes.h" @@ -23,8 +24,8 @@ // Cleanup the IR after manipulating it. passManager.addPass(createInlinerPass()); - passManager.addNestedPass<FuncOp>(createCanonicalizerPass()); - passManager.addNestedPass<FuncOp>(createCSEPass()); + passManager.addNestedPass<func::FuncOp>(createCanonicalizerPass()); + passManager.addNestedPass<func::FuncOp>(createCSEPass()); passManager.addPass(createSymbolDCEPass()); }
diff --git a/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp b/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp index ce174bc..c643cfc 100644 --- a/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp +++ b/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp
@@ -46,7 +46,7 @@ void runOnOperation() override { auto moduleOp = getOperation(); - SmallVector<FuncOp, 4> entryFuncOps; + SmallVector<func::FuncOp, 4> entryFuncOps; for (auto funcOp : moduleOp.getOps<FuncOp>()) { if (funcOp.isPublic() && !funcOp->hasAttr("iree.abi.stub")) { entryFuncOps.push_back(funcOp); @@ -103,13 +103,13 @@ // // NOTE: today we only support a single entry point; with minor tweaks we // could fix this up to support multiple if we wanted. - FuncOp createWrapperFunc(FuncOp entryFuncOp) { + func::FuncOp createWrapperFunc(FuncOp entryFuncOp) { // Convert argument types to those required by the binding ABI. // // NOTE: this is where we could change our signature to provide additional // values from the runtime bindings as may be required - like semaphores for // async behavior or cancellation. - auto entryFuncType = entryFuncOp.getType(); + auto entryFuncType = entryFuncOp.getFunctionType(); SmallVector<Type> inputTypes; for (auto oldType : entryFuncType.getInputs()) { inputTypes.push_back(mapToABIType(oldType));
diff --git a/iree/compiler/Bindings/TFLite/Transforms/Passes.cpp b/iree/compiler/Bindings/TFLite/Transforms/Passes.cpp index 6ddf1e5..088c4fe 100644 --- a/iree/compiler/Bindings/TFLite/Transforms/Passes.cpp +++ b/iree/compiler/Bindings/TFLite/Transforms/Passes.cpp
@@ -8,6 +8,7 @@ #include <memory> +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/PassOptions.h" #include "mlir/Pass/PassRegistry.h" #include "mlir/Transforms/Passes.h" @@ -23,8 +24,8 @@ // Cleanup the IR after manipulating it. passManager.addPass(createInlinerPass()); - passManager.addNestedPass<FuncOp>(createCanonicalizerPass()); - passManager.addNestedPass<FuncOp>(createCSEPass()); + passManager.addNestedPass<func::FuncOp>(createCanonicalizerPass()); + passManager.addNestedPass<func::FuncOp>(createCSEPass()); passManager.addPass(createSymbolDCEPass()); }
diff --git a/iree/compiler/Bindings/TFLite/Transforms/WrapEntryPoints.cpp b/iree/compiler/Bindings/TFLite/Transforms/WrapEntryPoints.cpp index e330c8d..ec94030 100644 --- a/iree/compiler/Bindings/TFLite/Transforms/WrapEntryPoints.cpp +++ b/iree/compiler/Bindings/TFLite/Transforms/WrapEntryPoints.cpp
@@ -131,8 +131,8 @@ // Creates dynamic dim globals for each input and output of |funcOp|. static std::pair<SmallVector<DynamicDims>, SmallVector<DynamicDims>> createDynamicDimGlobals(Location loc, StringRef namePrefix, - mlir::FuncOp funcOp, OpBuilder &moduleBuilder) { - auto funcType = funcOp.getType(); + mlir::func::FuncOp funcOp, OpBuilder &moduleBuilder) { + auto funcType = funcOp.getFunctionType(); // TFLite requires the tensor names at runtime. If they've previously been // extracted into iree.identifiers we use those and otherwise fallback to @@ -183,14 +183,14 @@ } // Derives a shape calculation function from the given entry point |funcOp|. - static mlir::FuncOp createShapeCalculationFunc( - Location loc, StringRef namePrefix, mlir::FuncOp funcOp, + static mlir::func::FuncOp createShapeCalculationFunc( + Location loc, StringRef namePrefix, mlir::func::FuncOp funcOp, ArrayRef<DynamicDims> inputDynamicDims, ArrayRef<DynamicDims> outputDynamicDims, IREE::Util::GlobalOp dirtyGlobalOp, OpBuilder &moduleBuilder) { // Clone the entire entry function with all its IR. auto calcFuncOp = - cast<mlir::FuncOp>(moduleBuilder.clone(*funcOp.getOperation())); + cast<mlir::func::FuncOp>(moduleBuilder.clone(*funcOp.getOperation())); calcFuncOp.setName( moduleBuilder.getStringAttr(namePrefix.str() + "_calculate_shapes")); calcFuncOp.setPrivate(); @@ -364,7 +364,7 @@ void createQueryInputShapeFunc(Location loc, StringRef namePrefix, ArrayRef<DynamicDims> inputDynamicDims, OpBuilder &moduleBuilder) { - auto queryFuncOp = moduleBuilder.create<mlir::FuncOp>( + auto queryFuncOp = moduleBuilder.create<mlir::func::FuncOp>( loc, namePrefix.str() + "_query_input_shape", moduleBuilder.getFunctionType(/*inputs=*/ TypeRange{ @@ -397,7 +397,7 @@ ArrayRef<DynamicDims> inputDynamicDims, IREE::Util::GlobalOp dirtyGlobalOp, OpBuilder &moduleBuilder) { - auto resizeFuncOp = moduleBuilder.create<mlir::FuncOp>( + auto resizeFuncOp = moduleBuilder.create<mlir::func::FuncOp>( loc, namePrefix.str() + "_resize_input_shape", moduleBuilder.getFunctionType(/*inputs=*/ TypeRange{ @@ -432,7 +432,7 @@ // func @_query_output_shape(%index : index, %shape : !util.list<index>) void createQueryOutputShapeFunc(Location loc, StringRef namePrefix, ArrayRef<DynamicDims> outputDynamicDims, - mlir::FuncOp calculateShapeFuncOp, + mlir::func::FuncOp calculateShapeFuncOp, OpBuilder &moduleBuilder) { auto queryFuncOp = moduleBuilder.create<FuncOp>( loc, namePrefix.str() + "_query_output_shape", @@ -474,7 +474,7 @@ // // NOTE: today we only support a single entry point; with minor tweaks we // could fix this up to support multiple if we wanted. - void createWrapperFunc(StringRef namePrefix, mlir::FuncOp entryFuncOp, + void createWrapperFunc(StringRef namePrefix, mlir::func::FuncOp entryFuncOp, ArrayRef<DynamicDims> inputDynamicDims, ArrayRef<DynamicDims> outputDynamicDims, IREE::Util::GlobalOp dirtyGlobalOp, @@ -482,14 +482,14 @@ // NOTE: this is where we could change our signature to provide additional // values from the runtime bindings as may be required - like semaphores for // async behavior or cancellation. - auto entryFuncType = entryFuncOp.getType(); + auto entryFuncType = entryFuncOp.getFunctionType(); auto bufferType = moduleBuilder.getType<IREE::HAL::BufferType>(); SmallVector<Type> inputTypes(entryFuncType.getNumInputs(), bufferType); SmallVector<Type> outputTypes(entryFuncType.getNumResults(), bufferType); auto wrapperFuncType = moduleBuilder.getFunctionType(inputTypes, outputTypes); - auto wrapperFuncOp = moduleBuilder.create<mlir::FuncOp>( + auto wrapperFuncOp = moduleBuilder.create<mlir::func::FuncOp>( entryFuncOp.getLoc(), "_tflite_main", wrapperFuncType); wrapperFuncOp.setPublic(); wrapperFuncOp.getOperation()->setAttr("iree.abi.stub", @@ -563,7 +563,7 @@ callResults); } - void wrapEntryPoint(mlir::FuncOp funcOp) { + void wrapEntryPoint(mlir::func::FuncOp funcOp) { auto loc = funcOp.getLoc(); auto namePrefix = ("_tflite_" + funcOp.getName()).str(); OpBuilder moduleBuilder(funcOp); @@ -606,8 +606,8 @@ // Populates attributes on |wrapperFuncOp| to support runtime reflection like // IO tensor names and quantization information. - void populateReflectionAttrs(mlir::FuncOp entryFuncOp, - mlir::FuncOp wrapperFuncOp) { + void populateReflectionAttrs(mlir::func::FuncOp entryFuncOp, + mlir::func::FuncOp wrapperFuncOp) { SmallVector<NamedAttribute, 4> attrs; attrs.push_back(buildIONamesAttr(entryFuncOp)); // TODO(#3972): tfl.io.quant: quantization information. @@ -620,7 +620,7 @@ // tfl.io.names=arg0;arg1;ret0;ret1 // // Default names will be used if no iree.identifiers are set on the function. - NamedAttribute buildIONamesAttr(mlir::FuncOp entryFuncOp) { + NamedAttribute buildIONamesAttr(mlir::func::FuncOp entryFuncOp) { SmallVector<std::string, 4> pieces; for (int i = 0; i < entryFuncOp.getNumArguments(); ++i) { auto identifierAttr =
diff --git a/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp b/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp index 17f9d5d..4b2d5de 100644 --- a/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp +++ b/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp
@@ -513,7 +513,7 @@ } } -LogicalResult createTensorEquivalenceClasses(FuncOp funcOp, +LogicalResult createTensorEquivalenceClasses(func::FuncOp funcOp, BufferizationPlan &plan) { auto bufferMappingFn = [&](Operation *op) -> WalkResult { return TypeSwitch<Operation *, LogicalResult>(op)
diff --git a/iree/compiler/Codegen/Common/BufferizationAnalysis.h b/iree/compiler/Codegen/Common/BufferizationAnalysis.h index e324eeb..aa936bf 100644 --- a/iree/compiler/Codegen/Common/BufferizationAnalysis.h +++ b/iree/compiler/Codegen/Common/BufferizationAnalysis.h
@@ -16,6 +16,7 @@ #include "llvm/ADT/EquivalenceClasses.h" #include "llvm/Support/Debug.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/Value.h" @@ -97,7 +98,7 @@ /// Analysis the `tensor` values in `funcOp` and groups them together into /// equivalence classes such that each class contains tensors that can be mapped /// to the same buffer. -LogicalResult createTensorEquivalenceClasses(FuncOp funcOp, +LogicalResult createTensorEquivalenceClasses(func::FuncOp funcOp, BufferizationPlan &plan); } // namespace iree_compiler
diff --git a/iree/compiler/Codegen/Common/DemoteF32ToF16.cpp b/iree/compiler/Codegen/Common/DemoteF32ToF16.cpp index 7e09857..0b586e5 100644 --- a/iree/compiler/Codegen/Common/DemoteF32ToF16.cpp +++ b/iree/compiler/Codegen/Common/DemoteF32ToF16.cpp
@@ -133,19 +133,19 @@ FloatTypeConverter converter; RewritePatternSet patterns(&getContext()); patterns.insert<GenericTypeConvert>(context, converter); - populateFunctionOpInterfaceTypeConversionPattern<FuncOp>(patterns, - converter); + populateFunctionOpInterfaceTypeConversionPattern<func::FuncOp>(patterns, + converter); ConversionTarget target(*context); // Operations are legal if they don't contain any illegal type. target.markUnknownOpDynamicallyLegal([](Operation *op) { if (auto globalOp = dyn_cast<IREE::Util::GlobalOp>(op)) { return !isIllegalType(globalOp.type()); } - if (auto funcOp = dyn_cast<FuncOp>(op)) { - for (Type type : funcOp.getType().getInputs()) { + if (auto funcOp = dyn_cast<func::FuncOp>(op)) { + for (Type type : funcOp.getFunctionType().getInputs()) { if (isIllegalType(type)) return false; } - for (Type type : funcOp.getType().getResults()) { + for (Type type : funcOp.getFunctionType().getResults()) { if (isIllegalType(type)) return false; } }
diff --git a/iree/compiler/Codegen/Common/DestructiveUpdateUtils.cpp b/iree/compiler/Codegen/Common/DestructiveUpdateUtils.cpp index 592ed8d..5bc2203 100644 --- a/iree/compiler/Codegen/Common/DestructiveUpdateUtils.cpp +++ b/iree/compiler/Codegen/Common/DestructiveUpdateUtils.cpp
@@ -294,11 +294,11 @@ // Return true if any control flow is found in the DispatchWorkgroupsOp besides // scf::ForOp. -static bool hasNonScfForControlFlow(FuncOp funcOp) { +static bool hasNonScfForControlFlow(func::FuncOp funcOp) { return funcOp ->walk([&](Operation *op) { if (isa<BranchOpInterface>(op) || isa<RegionBranchOpInterface>(op)) { - if (!isa<scf::ForOp, scf::IfOp>(op) && + if (!isa<scf::ForOp, scf::IfOp>(op) && !isa<linalg::LinalgOp>(op) && !isa<IREE::Flow::DispatchWorkgroupsOp>(op)) return WalkResult::interrupt(); } @@ -335,7 +335,7 @@ return success(); } -LogicalResult rewriteLinalgDestructiveUpdates(FuncOp funcOp) { +LogicalResult rewriteLinalgDestructiveUpdates(func::FuncOp funcOp) { // Bail on any control-flow for now. if (hasNonScfForControlFlow(funcOp)) return success();
diff --git a/iree/compiler/Codegen/Common/DestructiveUpdateUtils.h b/iree/compiler/Codegen/Common/DestructiveUpdateUtils.h index 04197ec..a0bdc7a 100644 --- a/iree/compiler/Codegen/Common/DestructiveUpdateUtils.h +++ b/iree/compiler/Codegen/Common/DestructiveUpdateUtils.h
@@ -7,6 +7,7 @@ #ifndef IREE_COMPILER_CODEGEN_COMMON_DESTRUCTIVEUPDATEUTILS_H_ #define IREE_COMPILER_CODEGEN_COMMON_DESTRUCTIVEUPDATEUTILS_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/BuiltinOps.h" namespace mlir { @@ -78,7 +79,7 @@ // tensor<...> -> !flow.dispatch.tensor<writeonly:...> // ``` // is elided. -LogicalResult rewriteLinalgDestructiveUpdates(FuncOp parentOp); +LogicalResult rewriteLinalgDestructiveUpdates(func::FuncOp parentOp); } // namespace iree_compiler } // namespace mlir
diff --git a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir index 5e40783..d8f5ae0 100644 --- a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir +++ b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
@@ -2206,7 +2206,7 @@ // ----- -builtin.func @tensor_insert_slice() { +func.func @tensor_insert_slice() { %c0 = arith.constant 0 : index %1 = hal.interface.constant.load[0] : index %2 = hal.interface.constant.load[1] : index @@ -2256,7 +2256,7 @@ // ----- -builtin.func @dynamic_update_slice() { +func.func @dynamic_update_slice() { %c0 = arith.constant 0 : index %c3 = arith.constant 3 : index %c0_i32 = arith.constant 0 : i32
diff --git a/iree/compiler/Codegen/Common/test/remove_trivial_loops.mlir b/iree/compiler/Codegen/Common/test/remove_trivial_loops.mlir index c5b6022..0505aae 100644 --- a/iree/compiler/Codegen/Common/test/remove_trivial_loops.mlir +++ b/iree/compiler/Codegen/Common/test/remove_trivial_loops.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-codegen-remove-single-iteration-loop))))' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(func.func(iree-codegen-remove-single-iteration-loop))))' %s | FileCheck %s #executable_layout = #hal.executable.layout<push_constants = 1, sets = [ #hal.descriptor_set.layout<0, bindings = [ @@ -14,7 +14,7 @@ workgroup_size = [64: index, 1: index, 1:index] } builtin.module { - builtin.func @dispatch_0() { + func.func @dispatch_0() { %c2 = arith.constant 2 : index %c256 = arith.constant 256 : index // CHECK: %[[C250:.+]] = arith.constant 250 : index @@ -61,7 +61,7 @@ translation_info = #translation } builtin.module { - builtin.func @workgroup_tile_loop() { + func.func @workgroup_tile_loop() { %c2048 = arith.constant 2048 : index %workgroup_id_x = hal.interface.workgroup.id[0] : index %workgroup_count_x = hal.interface.workgroup.count[0] : index @@ -95,7 +95,7 @@ translation_info = #translation } builtin.module { - builtin.func @workgroup_tile_loop_negative() { + func.func @workgroup_tile_loop_negative() { %c2048 = arith.constant 2048 : index %workgroup_id_x = hal.interface.workgroup.id[0] : index %workgroup_count_x = hal.interface.workgroup.count[0] : index @@ -132,7 +132,7 @@ workgroup_size = [8: index, 2: index, 1: index] } builtin.module { - builtin.func @both_workgroup_and_workitem() { + func.func @both_workgroup_and_workitem() { %c8 = arith.constant 8 : index %c32 = arith.constant 32 : index %c112 = arith.constant 112 : index
diff --git a/iree/compiler/Codegen/Common/test/tile_and_distribute_to_workgroups.mlir b/iree/compiler/Codegen/Common/test/tile_and_distribute_to_workgroups.mlir index 8aef8cb..2a9a331 100644 --- a/iree/compiler/Codegen/Common/test/tile_and_distribute_to_workgroups.mlir +++ b/iree/compiler/Codegen/Common/test/tile_and_distribute_to_workgroups.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-codegen-tile-and-distribute-to-workgroups)))), canonicalize, cse' -split-input-file %s | FileCheck %s +// RUN: iree-opt -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(func.func(iree-codegen-tile-and-distribute-to-workgroups)))), canonicalize, cse' -split-input-file %s | FileCheck %s #config = #iree_codegen.lowering_config<tile_sizes = [[64, 64, 0], [16, 4, 64], [4, 4, 4]], native_vector_size = [4, 4, 4]> #executable_layout = #hal.executable.layout<push_constants = 0, sets = [
diff --git a/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir b/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir index 06613df..b8042ef 100644 --- a/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir +++ b/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir
@@ -11,7 +11,7 @@ // ----- module { - builtin.func @test() attributes { + func.func @test() attributes { translation_info = #iree_codegen.translation_info<CPUDefault>} { return } @@ -21,7 +21,7 @@ // ----- module { - builtin.func @test() attributes { + func.func @test() attributes { lowering_config = #iree_codegen.lowering_config<tile_sizes = []>} { return }
diff --git a/iree/compiler/Codegen/Interfaces/BufferizationInterfaces.cpp b/iree/compiler/Codegen/Interfaces/BufferizationInterfaces.cpp index 9302dd1..1818ddb 100644 --- a/iree/compiler/Codegen/Interfaces/BufferizationInterfaces.cpp +++ b/iree/compiler/Codegen/Interfaces/BufferizationInterfaces.cpp
@@ -279,10 +279,13 @@ vector::registerBufferizableOpInterfaceExternalModels(registry); // Register IREE operations. - registry.addOpInterface<IREE::Flow::DispatchTensorLoadOp, - DispatchTensorLoadOpInterface>(); - registry.addOpInterface<IREE::Flow::DispatchTensorStoreOp, - DispatchTensorStoreOpInterface>(); + registry.addExtension( + +[](MLIRContext *ctx, IREE::Flow::FlowDialect *dialect) { + IREE::Flow::DispatchTensorLoadOp::attachInterface< + DispatchTensorLoadOpInterface>(*ctx); + IREE::Flow::DispatchTensorStoreOp::attachInterface< + DispatchTensorStoreOpInterface>(*ctx); + }); } void addPostAnalysisTransformations(OneShotBufferizationOptions &options) {
diff --git a/iree/compiler/Codegen/Interfaces/ProcessorOpInterfaces.cpp b/iree/compiler/Codegen/Interfaces/ProcessorOpInterfaces.cpp index 9fe9962..2fabe05 100644 --- a/iree/compiler/Codegen/Interfaces/ProcessorOpInterfaces.cpp +++ b/iree/compiler/Codegen/Interfaces/ProcessorOpInterfaces.cpp
@@ -6,6 +6,7 @@ #include "iree/compiler/Codegen/Interfaces/ProcessorOpInterfaces.h" +#include "iree/compiler/Dialect/HAL/IR/HALDialect.h" #include "iree/compiler/Dialect/HAL/IR/HALOps.h" #include "mlir/Dialect/GPU/GPUDialect.h" @@ -75,15 +76,19 @@ }; void registerProcessorOpInterfaceExternalModels(DialectRegistry ®istry) { - registry.addOpInterface<gpu::ThreadIdOp, ThreadIdOpInterface>(); - registry.addOpInterface<gpu::BlockDimOp, BlockDimOpInterface>(); + registry.addExtension(+[](MLIRContext *ctx, gpu::GPUDialect *dialect) { + gpu::ThreadIdOp::attachInterface<ThreadIdOpInterface>(*ctx); + gpu::BlockDimOp::attachInterface<BlockDimOpInterface>(*ctx); + }); - registry.addOpInterface<IREE::HAL::InterfaceWorkgroupIDOp, - WorkgroupIdOpInterface>(); - registry.addOpInterface<IREE::HAL::InterfaceWorkgroupCountOp, - WorkgroupCountOpInterface>(); - registry.addOpInterface<IREE::HAL::InterfaceWorkgroupSizeOp, - WorkgroupTileSizeOpInterface>(); + registry.addExtension(+[](MLIRContext *ctx, IREE::HAL::HALDialect *dialect) { + IREE::HAL::InterfaceWorkgroupIDOp::attachInterface<WorkgroupIdOpInterface>( + *ctx); + IREE::HAL::InterfaceWorkgroupCountOp::attachInterface< + WorkgroupCountOpInterface>(*ctx); + IREE::HAL::InterfaceWorkgroupSizeOp::attachInterface< + WorkgroupTileSizeOpInterface>(*ctx); + }); } } // namespace iree_compiler
diff --git a/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp b/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp index 12963e1..b2f5db9 100644 --- a/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp +++ b/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
@@ -627,7 +627,7 @@ public: explicit ConvertHALEntryPointFuncOp(MLIRContext *context, LLVMTypeConverter &converter) - : ConvertToLLVMPattern(mlir::FuncOp::getOperationName(), context, + : ConvertToLLVMPattern(mlir::func::FuncOp::getOperationName(), context, converter, 100) {} LogicalResult matchAndRewrite( @@ -635,7 +635,7 @@ ConversionPatternRewriter &rewriter) const override { auto stdFuncOp = cast<FuncOp>(op); if (!stdFuncOp.isPublic()) return failure(); - FunctionType fnType = stdFuncOp.getType(); + FunctionType fnType = stdFuncOp.getFunctionType(); if (fnType.getNumInputs() != 0 || fnType.getNumResults() != 0) { op->emitWarning() << "public functions on executables must be () -> ()"; return failure();
diff --git a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir index 4f219c4..56bc1b3 100644 --- a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir +++ b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
@@ -227,7 +227,7 @@ hal.executable.variant @system_elf_x86_64, target = <"llvm", "system-elf-x86_64"> { hal.executable.entry_point @preset_config layout(#executable_layout) builtin.module { - builtin.func @preset_config() { + func.func @preset_config() { %cst = arith.constant 0.000000e+00 : f32 %lhs_binding = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:128x256xf32> @@ -272,7 +272,7 @@ hal.executable.variant @system_elf_x86_64, target = <"llvm", "system-elf-x86_64"> { hal.executable.entry_point @tensor_insert_slice layout(#executable_layout) builtin.module { - builtin.func @tensor_insert_slice() { + func.func @tensor_insert_slice() { %d0 = hal.interface.constant.load[0] : index %d1 = hal.interface.constant.load[1] : index %d2 = hal.interface.constant.load[2] : index @@ -315,7 +315,7 @@ hal.executable.variant @system_elf_x86_64, target = <"llvm", "system-elf-x86_64"> { hal.executable.entry_point @extract_slice layout(#executable_layout) builtin.module { - builtin.func @extract_slice() { + func.func @extract_slice() { %d0 = hal.interface.constant.load[0] : index %d1 = hal.interface.constant.load[1] : index %d2 = hal.interface.constant.load[2] : index @@ -359,7 +359,7 @@ hal.executable.variant @system_elf_x86_64, target = <"llvm", "system-elf-x86_64"> { hal.executable.entry_point @static_1d_fft_stage2 layout(#executable_layout) builtin.module { - builtin.func @static_1d_fft_stage2() { + func.func @static_1d_fft_stage2() { %c0 = arith.constant 0 : index %c2 = arith.constant 2 : index %cst = arith.constant dense<[1.000000e+00, 6.12323426E-17]> : tensor<2xf32> @@ -397,7 +397,7 @@ hal.executable.variant @system_elf_x86_64, target = <"llvm", "system-elf-x86_64"> { hal.executable.entry_point @static_3d_fft_stage3 layout(#executable_layout) builtin.module { - builtin.func @static_3d_fft_stage3() { + func.func @static_3d_fft_stage3() { %c3 = arith.constant 3 : index %cst = arith.constant dense<[1.000000e+00, 0.707106769, 6.12323426E-17, -0.707106769]> : tensor<4xf32> %cst_0 = arith.constant dense<[-0.000000e+00, -0.707106769, -1.000000e+00, -0.707106769]> : tensor<4xf32> @@ -435,7 +435,7 @@ hal.executable.variant @system_elf_x86_64, target = <"llvm", "system-elf-x86_64"> { hal.executable.entry_point @outs_fusion_fn layout(#executable_layout) builtin.module { - builtin.func @outs_fusion_fn() { + func.func @outs_fusion_fn() { %cst = arith.constant 0.0 : f32 %d0 = hal.interface.constant.load[0] : index %d1 = hal.interface.constant.load[1] : index
diff --git a/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp b/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp index 4cce8cf..2ebdf28 100644 --- a/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp +++ b/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
@@ -210,13 +210,13 @@ class ConvertFunc : public ConvertToLLVMPattern { public: explicit ConvertFunc(MLIRContext *context, LLVMTypeConverter &converter) - : ConvertToLLVMPattern(mlir::FuncOp::getOperationName(), context, + : ConvertToLLVMPattern(mlir::func::FuncOp::getOperationName(), context, converter, 100) {} LogicalResult matchAndRewrite( Operation *op, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { - auto funcOp = cast<FuncOp>(op); - FunctionType fnType = funcOp.getType(); + auto funcOp = cast<func::FuncOp>(op); + FunctionType fnType = funcOp.getFunctionType(); (void)fnType; if (!funcOp.isPublic()) return failure();
diff --git a/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp index 88b8457..69795a1 100644 --- a/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp +++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp
@@ -128,7 +128,8 @@ } /// Return a flattened Id Value by combining the 3D gpu thread IDs. -static Value createFlatId(FuncOp funcOp, ArrayRef<int64_t> workgroupSize) { +static Value createFlatId(func::FuncOp funcOp, + ArrayRef<int64_t> workgroupSize) { OpBuilder b(funcOp.getBody()); Type indexType = b.getIndexType(); AffineExpr d0 = getAffineDimExpr(0, b.getContext()); @@ -148,7 +149,7 @@ } /// Distribute a transfer read operations on the given thread ids. -static void distributeTransferRead(FuncOp funcOp, Value flatThreadId, +static void distributeTransferRead(func::FuncOp funcOp, Value flatThreadId, int64_t flatWorkgroupSize) { funcOp.walk([&](vector::TransferReadOp readOp) { OpBuilder b(readOp);
diff --git a/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.cpp index 0963c3b..c7ab1aa 100644 --- a/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.cpp +++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.cpp
@@ -52,7 +52,7 @@ return procInfo; } -std::array<int64_t, 3> getWorkgroupSize(mlir::FuncOp funcOp) { +std::array<int64_t, 3> getWorkgroupSize(mlir::func::FuncOp funcOp) { std::array<int64_t, 3> workgroupSize; auto entryPointOp = mlir::iree_compiler::getEntryPoint(funcOp); llvm::Optional<mlir::ArrayAttr> workgroupSizeAttr =
diff --git a/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.h b/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.h index 35843f6..aaca6bd 100644 --- a/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.h +++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.h
@@ -30,7 +30,7 @@ llvm::ArrayRef<int64_t> numSubgroups); /// return the workgroup size associated to the funcOp entry point. -std::array<int64_t, 3> getWorkgroupSize(mlir::FuncOp funcOp); +std::array<int64_t, 3> getWorkgroupSize(mlir::func::FuncOp funcOp); } // namespace iree_compiler } // namespace mlir
diff --git a/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir index 7004ce8..6120f0a 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-llvmgpu-tile-and-distribute))))' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(func.func(iree-llvmgpu-tile-and-distribute))))' %s | FileCheck %s #config = #iree_codegen.lowering_config<tile_sizes = [[2, 256, 4]]> #translation = #iree_codegen.translation_info<LLVMGPUMatmulSimt, workload_per_wg = [256, 2]> @@ -22,7 +22,7 @@ workgroup_size = [64 : index, 1 : index, 1 : index] } builtin.module { - builtin.func @dot_dispatch_0() { + func.func @dot_dispatch_0() { %cst = arith.constant 0.000000e+00 : f32 %c0 = arith.constant 0 : index %c1024 = arith.constant 1024 : index @@ -181,7 +181,7 @@ workgroup_size = [64 : index, 8 : index, 1 : index] } builtin.module { - builtin.func @dot_dispatch_0() { + func.func @dot_dispatch_0() { %cst = arith.constant 0.000000e+00 : f32 %c0 = arith.constant 0 : index %c1024 = arith.constant 1024 : index @@ -263,7 +263,7 @@ workgroup_size = [1: index, 1: index, 1: index] } builtin.module { - builtin.func @predict_dispatch_153() { + func.func @predict_dispatch_153() { %c0 = arith.constant 0 : index %cst = arith.constant 0x7FC00000 : f32 %cst_0 = arith.constant 0xFF800000 : f32
diff --git a/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir b/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir index 5511cce..dde4b51 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-llvmgpu-distribute-shared-memory-copy))))' -cse %s | FileCheck %s +// RUN: iree-opt -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(func.func(iree-llvmgpu-distribute-shared-memory-copy))))' -cse %s | FileCheck %s // CHECK-DAG: #[[$MAP0:.*]] = affine_map<()[s0, s1, s2] -> (s1 * 8 + s2 * 32 + s0 floordiv 4)> // CHECK-DAG: #[[$MAP1:.*]] = affine_map<()[s0] -> (s0 * 4 - (s0 floordiv 4) * 16)> @@ -25,7 +25,7 @@ memref.global "private" @__shared_memory___0 : memref<256x4xf32, 3> memref.global "private" @__shared_memory__ : memref<64x16xf32, 3> // CHECK-LABEL: @shared_mem_cpy( - builtin.func @shared_mem_cpy( + func.func @shared_mem_cpy( %m0 : memref<64x16xf32>, %m1 : memref<256x4xf32>, %m2 : memref<3x512xf32>) { %c0 = arith.constant 0 : index
diff --git a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir index 25370ff..68af7ed 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
@@ -134,7 +134,7 @@ hal.executable.variant @cuda, target = <"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @tensor_insert_slice layout(#executable_layout) builtin.module { - builtin.func @tensor_insert_slice() { + func.func @tensor_insert_slice() { %c0 = arith.constant 0 : index %size_y = hal.interface.constant.load[0] : index %size_x = hal.interface.constant.load[1] : index @@ -178,7 +178,7 @@ hal.executable.variant @cuda, target = <"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @copy_as_generic layout(#executable_layout) builtin.module { - builtin.func @copy_as_generic() { + func.func @copy_as_generic() { %c0 = arith.constant 0 : index %d0 = hal.interface.constant.load[0] : index %d1 = hal.interface.constant.load[1] : index @@ -214,7 +214,7 @@ hal.executable.variant @cuda, target = <"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @static_1d_fft_stage2 layout(#executable_layout) builtin.module { - builtin.func @static_1d_fft_stage2() { + func.func @static_1d_fft_stage2() { %c0 = arith.constant 0 : index %c2 = arith.constant 2 : index %cst = arith.constant dense<[1.000000e+00, 6.12323426E-17]> : tensor<2xf32> @@ -251,7 +251,7 @@ hal.executable.variant @cuda, target = <"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @static_3d_fft_stage3 layout(#executable_layout) builtin.module { - builtin.func @static_3d_fft_stage3() { + func.func @static_3d_fft_stage3() { %c0 = arith.constant 0 : index %c3 = arith.constant 3 : index %c64 = arith.constant 64 : index
diff --git a/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir b/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir index 0e2cf86..06444d8 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
@@ -294,7 +294,7 @@ hal.executable.variant @cuda, target = <"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @vector_add_dispatch layout(#executable_layout) builtin.module { - builtin.func @vector_add_dispatch() { + func.func @vector_add_dispatch() { %c0 = arith.constant 0 : index %c16384 = arith.constant 16384 : index %0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:16384xf32> @@ -343,7 +343,7 @@ hal.executable.variant @cuda, target = <"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @vector_reduction_dispatch layout(#executable_layout) builtin.module { - builtin.func @vector_reduction_dispatch() { + func.func @vector_reduction_dispatch() { %c0 = arith.constant 0 : index %c16384 = arith.constant 16384 : index %cst = arith.constant 1.000000e+00 : f32
diff --git a/iree/compiler/Codegen/PassDetail.h b/iree/compiler/Codegen/PassDetail.h index 4ac9280..501417a 100644 --- a/iree/compiler/Codegen/PassDetail.h +++ b/iree/compiler/Codegen/PassDetail.h
@@ -8,6 +8,7 @@ #define IREE_COMPILER_CONVERSIONS_PASS_DETAIL_H_ #include "iree/compiler/Dialect/HAL/IR/HALOps.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir {
diff --git a/iree/compiler/Codegen/Passes.td b/iree/compiler/Codegen/Passes.td index f5d20a4..9637a2f 100644 --- a/iree/compiler/Codegen/Passes.td +++ b/iree/compiler/Codegen/Passes.td
@@ -14,14 +14,14 @@ //------------------------------------------------------------------------------ def CleanupBufferAllocView : - Pass<"iree-codegen-cleanup-buffer-alloc-view", "FuncOp"> { + Pass<"iree-codegen-cleanup-buffer-alloc-view", "func::FuncOp"> { let summary = "Performs cleanups over HAL interface/buffer allocation/view operations"; let constructor = "mlir::iree_compiler::createCleanupBufferAllocViewPass()"; } def ConvertToDestinationPassingStyle : - Pass<"iree-codegen-convert-to-destination-passing-style", "FuncOp"> { + Pass<"iree-codegen-convert-to-destination-passing-style", "func::FuncOp"> { let summary = "Transforms the code to make the dispatch use destination-passing style"; let constructor = "mlir::iree_compiler::createConvertToDestinationPassingStylePass()"; @@ -41,7 +41,7 @@ } def FoldAffineMinInDistributedLoops : - Pass<"iree-codegen-fold-affinemin-in-distributed-loops", "FuncOp"> { + Pass<"iree-codegen-fold-affinemin-in-distributed-loops", "func::FuncOp"> { let summary = "Fold `affine.min` ops in distributed loops"; let constructor = "mlir::iree_compiler::createFoldAffineMinInDistributedLoopsPass()"; } @@ -53,14 +53,14 @@ } def ForOpCanonicalization : - Pass<"iree-codegen-canonicalize-scf-for", "FuncOp"> { + Pass<"iree-codegen-canonicalize-scf-for", "func::FuncOp"> { let summary = "Adhoc canonicalization of selected loop-carried values/dependencies for scf.for ops"; let constructor = "mlir::iree_compiler::createForOpCanonicalizationPass()"; } def LinalgBufferize : - Pass<"iree-codegen-linalg-bufferize", "FuncOp"> { + Pass<"iree-codegen-linalg-bufferize", "func::FuncOp"> { let summary = "Convert from to Linalg ops on tensors to buffers"; let constructor = "mlir::iree_compiler::createLinalgBufferizePass(nullptr)"; } @@ -80,26 +80,26 @@ } def OptimizeVectorTransfer : - Pass<"iree-codegen-optimize-vector-transfer", "FuncOp"> { + Pass<"iree-codegen-optimize-vector-transfer", "func::FuncOp"> { let summary = "Run optimization transformations on vector transfer operations"; let constructor = "mlir::iree_compiler::createOptimizeVectorTransferPass()"; } def TileAndDistributeToWorkgroups : - Pass<"iree-codegen-tile-and-distribute-to-workgroups", "FuncOp"> { + Pass<"iree-codegen-tile-and-distribute-to-workgroups", "func::FuncOp"> { let summary = "Tile and distribute operations to workgroups"; let constructor = "mlir::iree_compiler::createTileAndDistributeToWorkgroupsPass()"; } def TypePropagation : - Pass<"iree-codegen-type-propagation", "FuncOp"> { + Pass<"iree-codegen-type-propagation", "func::FuncOp"> { let summary = "Propogate the type of tensor to avoid load/stores of illegal bit widths"; let constructor = "mlir::iree_compiler::createTypePropagationPass()"; } def RemoveSingleIterationLoop : - Pass<"iree-codegen-remove-single-iteration-loop", "FuncOp"> { + Pass<"iree-codegen-remove-single-iteration-loop", "func::FuncOp"> { let summary = "Remove distributed loop with single iteration."; let constructor = "mlir::iree_compiler::createRemoveSingleIterationLoopPass()"; } @@ -116,21 +116,21 @@ // TODO: Rename argument to be fully qualified. def LinalgToVectorVectorizeConv : - Pass<"iree-codegen-vectorize-linalg-conv", "FuncOp"> { + Pass<"iree-codegen-vectorize-linalg-conv", "func::FuncOp"> { let summary = "Vectorize a very specific form of linalg.conv"; let constructor = "mlir::iree_compiler::createLinalgToVectorVectorizeConvPass()"; } def LinalgToVectorVectorizeMMT4d : - Pass<"iree-codegen-vectorize-linalg-mmt4d", "FuncOp"> { + Pass<"iree-codegen-vectorize-linalg-mmt4d", "func::FuncOp"> { let summary = "Vectorize a very specific form of linalg.mmt4d"; let constructor = "mlir::iree_compiler::createLinalgToVectorVectorizeMMT4dPass()"; } def TensorToVectorVectorizePad : - Pass<"iree-codegen-vectorize-tensor-pad", "FuncOp"> { + Pass<"iree-codegen-vectorize-tensor-pad", "func::FuncOp"> { let summary = "Vectorize a very specific form of tensor.pad with " "control flows"; let constructor = @@ -145,7 +145,7 @@ } def MemrefCopyToLinalgPass : - Pass<"iree-codegen-memrefcopy-to-linalg", "FuncOp"> { + Pass<"iree-codegen-memrefcopy-to-linalg", "func::FuncOp"> { let summary = "Convert memref.copy to linalg op"; let constructor = "mlir::iree_compiler::createMemrefCopyToLinalgPass()"; @@ -183,20 +183,20 @@ } def LLVMCPUTileFuseAndVectorize - : Pass<"iree-llvmcpu-tile-fuse-and-vectorize", "FuncOp"> { + : Pass<"iree-llvmcpu-tile-fuse-and-vectorize", "func::FuncOp"> { let summary = "Tile, fuse and vectorize Linalg ops"; let constructor = "mlir::iree_compiler::createLLVMCPUTileFuseAndVectorizePass()"; } def LLVMCPUUnfuseFMAOps : - Pass<"iree-llvmcpu-unfuse-fma-pass", "FuncOp"> { + Pass<"iree-llvmcpu-unfuse-fma-pass", "func::FuncOp"> { let summary = "Convert llvm.fma into unfused mulf and addf ops"; let constructor = "mlir::iree_compiler::createLLVMCPUUnfuseFMAOpsPass()"; } def VectorContractCustomKernels : - Pass<"iree-llvmcpu-vector-contract-custom-kernels", "FuncOp"> { + Pass<"iree-llvmcpu-vector-contract-custom-kernels", "func::FuncOp"> { let summary = "Enable custom kernels (inline assembly or intrinsics) for some vector.contract ops"; let constructor = "mlir::iree_compiler::createVectorContractCustomKernelsPass()"; let options = [ @@ -240,49 +240,49 @@ } def LLVMGPUTileAndDistribute : - Pass<"iree-llvmgpu-tile-and-distribute", "FuncOp"> { + Pass<"iree-llvmgpu-tile-and-distribute", "func::FuncOp"> { let summary = "Pass to tile and distribute linalg ops within a workgroup."; let constructor = "mlir::iree_compiler::createLLVMGPUTileAndDistribute()"; } def LLVMGPUVectorization : - Pass<"iree-llvmgpu-vectorization", "FuncOp"> { + Pass<"iree-llvmgpu-vectorization", "func::FuncOp"> { let summary = "Pass to convert linalg into Vector."; let constructor = "mlir::iree_compiler::createLLVMGPUVectorizationPass()"; } def LLVMGPUTensorCoreVectorization : - Pass<"iree-llvmgpu-tensorcore-vectorization", "FuncOp"> { + Pass<"iree-llvmgpu-tensorcore-vectorization", "func::FuncOp"> { let summary = "Pass to convert linalg into Vector and transform it to a form that can be lowered to GPU MMA ops"; let constructor = "mlir::iree_compiler::createLLVMGPUTensorCoreVectorizationPass()"; } def LLVMGPUVectorLowering : - Pass<"iree-llvmgpu-vector-lowering", "FuncOp"> { + Pass<"iree-llvmgpu-vector-lowering", "func::FuncOp"> { let summary = "Pass to lower Vector ops before conversion to LLVM."; let constructor = "mlir::iree_compiler::createLLVMGPUVectorLoweringPass()"; } def LLVMGPUDistributeSharedMemoryCopy : - Pass<"iree-llvmgpu-distribute-shared-memory-copy", "FuncOp"> { + Pass<"iree-llvmgpu-distribute-shared-memory-copy", "func::FuncOp"> { let summary = "Pass to distribute shared memory copies to threads."; let constructor = "mlir::iree_compiler::createLLVMGPUDistributeSharedMemoryCopy()"; } def LLVMGPUPipelining : - Pass<"iree-llvmgpu-pipelining", "FuncOp"> { + Pass<"iree-llvmgpu-pipelining", "func::FuncOp"> { let summary = "Pass to do software pipelining."; let constructor = "mlir::iree_compiler::createLLVMGPUPipeliningPass()"; } def LLVMGPUMultiBuffering : - Pass<"iree-llvmgpu-multi-buffering", "FuncOp"> { + Pass<"iree-llvmgpu-multi-buffering", "func::FuncOp"> { let summary = "Pass to do multi buffering."; let constructor = "mlir::iree_compiler::createLLVMGPUMultiBuffering()"; } def LLVMGPUVectorToGPU : - Pass<"iree-llvmgpu-vector-to-gpu", "FuncOp"> { + Pass<"iree-llvmgpu-vector-to-gpu", "func::FuncOp"> { let summary = "Pass to convert vector to gpu."; let constructor = "mlir::iree_compiler::createLLVMGPUVectorToGPU()"; } @@ -311,24 +311,24 @@ let constructor = "mlir::iree_compiler::createSPIRVInitConfigPass()"; } -def SPIRVTile : Pass<"iree-spirv-tile", "FuncOp"> { +def SPIRVTile : Pass<"iree-spirv-tile", "func::FuncOp"> { let summary = "Tile Linalg ops with tensor semantics to invocations"; let constructor = "mlir::iree_compiler::createSPIRVTilePass()"; } -def SPIRVDistribute : Pass<"iree-spirv-distribute", "FuncOp"> { +def SPIRVDistribute : Pass<"iree-spirv-distribute", "func::FuncOp"> { let summary = "Distribute tiled loop nests to invocations"; let constructor = "mlir::iree_compiler::createSPIRVDistributePass()"; } -def SPIRVTileAndDistribute : Pass<"iree-spirv-tile-and-distribute", "FuncOp"> { +def SPIRVTileAndDistribute : Pass<"iree-spirv-tile-and-distribute", "func::FuncOp"> { let summary = "Tile and distribute Linalg ops with buffer semantics to " "invocations"; let constructor = "mlir::iree_compiler::createSPIRVTileAndDistributePass()"; } def SPIRVTileAndVectorizeToCooperativeOps : Pass< - "iree-spirv-tile-and-vectorize-to-cooperative-ops", "FuncOp"> { + "iree-spirv-tile-and-vectorize-to-cooperative-ops", "func::FuncOp"> { let summary = "Tile Linalg ops with buffer semantics to subgroups and " "vectorize to vector ops suitable for lowering to SPIR-V " "cooperative ops"; @@ -336,13 +336,13 @@ "mlir::iree_compiler::createSPIRVTileAndVectorizeToCooperativeOpsPass()"; } -def SPIRVVectorize : Pass<"iree-spirv-vectorize", "FuncOp"> { +def SPIRVVectorize : Pass<"iree-spirv-vectorize", "func::FuncOp"> { let summary = "Vectorize Linalg ops with buffer semantics"; let constructor = "mlir::iree_compiler::createSPIRVVectorizePass()"; } def SPIRVVectorToCooperativeOps : - Pass<"iree-spirv-vector-to-cooperative-ops", "FuncOp"> { + Pass<"iree-spirv-vector-to-cooperative-ops", "func::FuncOp"> { let summary = "Convert vector ops to SPIR-V cooperative ops"; let constructor = "mlir::iree_compiler::createSPIRVVectorToCooperativeOpsPass()"; @@ -355,13 +355,13 @@ } def SPIRVFuseTensorPadWithConsumer : - Pass<"iree-spirv-fuse-tensor-pad-with-consumer", "FuncOp"> { + Pass<"iree-spirv-fuse-tensor-pad-with-consumer", "func::FuncOp"> { let summary = "Fuse tensor.pad op into its consumer op's tiled loop nest"; let constructor = "mlir::iree_compiler::createSPIRVFuseTensorPadWithConsumerPass()"; } def SPIRVCreateFastSlowPath : - Pass<"iree-spirv-create-fast-slow-path", "FuncOp"> { + Pass<"iree-spirv-create-fast-slow-path", "func::FuncOp"> { let summary = "Create separate fast and slow paths to handle padding"; let constructor = "mlir::iree_compiler::createSPIRVCreateFastSlowPathPass()"; }
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp b/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp index d0fb79f..835760a 100644 --- a/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp +++ b/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp
@@ -489,7 +489,7 @@ FuncOp funcOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const { TypeConverter::SignatureConversion signatureConverter( - funcOp.getType().getNumInputs()); + funcOp.getFunctionType().getNumInputs()); for (const auto &arg : llvm::enumerate(funcOp.getArguments())) { if (memrefUsageAnalysis.shouldVectorizeMemRef(arg.value())) { if (auto memrefType = getVectorizedMemRefType(rewriter, arg.value())) {
diff --git a/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir b/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir index 005d010..7a491c2 100644 --- a/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir +++ b/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir
@@ -15,7 +15,7 @@ }> { hal.executable.entry_point @static_1d_sort layout(#executable_layout) builtin.module { - builtin.func @static_1d_sort() { + func.func @static_1d_sort() { %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readwrite:1000xi32> %1 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [1000], strides = [1] : !flow.dispatch.tensor<readwrite:1000xi32> -> tensor<1000xi32> @@ -60,7 +60,7 @@ }> { hal.executable.entry_point @static_3d_sort layout(#executable_layout) builtin.module { - builtin.func @static_3d_sort() { + func.func @static_3d_sort() { %c64 = arith.constant 64 : index %c128 = arith.constant 128 : index %c0 = arith.constant 0 : index @@ -109,7 +109,7 @@ }> { hal.executable.entry_point @static_1d_fft_stage2 layout(#executable_layout) builtin.module { - builtin.func @static_1d_fft_stage2() { + func.func @static_1d_fft_stage2() { %c0 = arith.constant 0 : index %c2 = arith.constant 2 : index %cst = arith.constant dense<[1.000000e+00, 6.12323426E-17]> : tensor<2xf32> @@ -154,7 +154,7 @@ }> { hal.executable.entry_point @static_3d_fft_stage3 layout(#executable_layout) builtin.module { - builtin.func @static_3d_fft_stage3() { + func.func @static_3d_fft_stage3() { %c0 = arith.constant 0 : index %c3 = arith.constant 3 : index %c64 = arith.constant 64 : index @@ -203,7 +203,7 @@ }> { hal.executable.entry_point @tensor_insert layout(#executable_layout) builtin.module { - builtin.func @tensor_insert() { + func.func @tensor_insert() { %offset_y = hal.interface.constant.load[0] : index %offset_x = hal.interface.constant.load[1] : index %source_size_y = hal.interface.constant.load[2] : index @@ -251,7 +251,7 @@ }> { hal.executable.entry_point @tensor_extract layout(#executable_layout) builtin.module { - builtin.func @tensor_extract() { + func.func @tensor_extract() { %offset_y = hal.interface.constant.load[0] : index %offset_x = hal.interface.constant.load[1] : index %source_size_y = hal.interface.constant.load[2] : index
diff --git a/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ops.mlir b/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ops.mlir index dd29936..a850da3 100644 --- a/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ops.mlir +++ b/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ops.mlir
@@ -16,7 +16,7 @@ }> { hal.executable.entry_point @copy_as_generic layout(#executable_layout) builtin.module { - builtin.func @copy_as_generic() { + func.func @copy_as_generic() { %c0 = arith.constant 0 : index %d0 = hal.interface.constant.load[0] : index %d1 = hal.interface.constant.load[1] : index @@ -60,7 +60,7 @@ }> { hal.executable.entry_point @copy layout(#executable_layout) builtin.module { - builtin.func @copy() { + func.func @copy() { %c0 = arith.constant 0 : index %c224 = arith.constant 224 : index %c3 = arith.constant 3 : index
diff --git a/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir b/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir index 4a9b306..c0a88e1 100644 --- a/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir +++ b/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir
@@ -1,5 +1,5 @@ // TODO(antiagainst): Fix promotion to workgroup and enable the test. -// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-tile-and-distribute,iree-spirv-vectorize,canonicalize,cse))))' | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(func.func(iree-spirv-tile-and-distribute,iree-spirv-vectorize,canonicalize,cse))))' | FileCheck %s hal.executable private @matmul_promote_workgroup_memory { hal.interface @io {
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute.mlir index 530531f..23663d7 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-tile-and-distribute))))' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(func.func(iree-spirv-tile-and-distribute))))' %s | FileCheck %s #map0 = affine_map<()[s0] -> (s0 * 8)> #map1 = affine_map<()[s0, s1] -> (8, s1 - s0 * 8)>
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir index e657a16..e214fad 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-tile-and-distribute))))' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(func.func(iree-spirv-tile-and-distribute))))' %s | FileCheck %s #config = #iree_codegen.lowering_config<tile_sizes = [[1, 16], [1, 1]]> #translation = #iree_codegen.translation_info<SPIRVDistribute, workload_per_wg = [16, 1]> @@ -16,7 +16,7 @@ workgroup_size = [16 : index, 1 : index, 1 : index] } builtin.module { - builtin.func @static_scatter_update_slice() { + func.func @static_scatter_update_slice() { %c40 = arith.constant 40 : index %c500 = arith.constant 500 : index %c0 = arith.constant 0 : index
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir index 36a3daf..d9279e6 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-tile-and-distribute, cse))))' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(func.func(iree-spirv-tile-and-distribute, cse))))' %s | FileCheck %s #config = #iree_codegen.lowering_config<tile_sizes = [[1, 0, 16], [1, 0, 1]]> #translation = #iree_codegen.translation_info<SPIRVDistribute, workload_per_wg = [16, 1]> @@ -15,7 +15,7 @@ workgroup_size = [16 : index, 1 : index, 1 : index] } builtin.module { - builtin.func @static_3d_sort() { + func.func @static_3d_sort() { %c64 = arith.constant 64 : index %c128 = arith.constant 128 : index %c0 = arith.constant 0 : index
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir index aaabd0c..f3f184d 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-set-num-workgroups,builtin.module(builtin.func(iree-spirv-tile,iree-spirv-vectorize))))' -cse %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-set-num-workgroups,builtin.module(func.func(iree-spirv-tile,iree-spirv-vectorize))))' -cse %s | FileCheck %s #config = #iree_codegen.lowering_config<tile_sizes = [[1, 8, 64], [1, 8, 4], [0, 0, 0, 4]]> #translation = #iree_codegen.translation_info<SPIRVVectorize, workload_per_wg = [64, 8, 1]>
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir index 631ce5e..245847c 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-set-num-workgroups,builtin.module(builtin.func(iree-spirv-create-fast-slow-path,iree-spirv-tile,iree-spirv-vectorize))))' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-set-num-workgroups,builtin.module(func.func(iree-spirv-create-fast-slow-path,iree-spirv-tile,iree-spirv-vectorize))))' %s | FileCheck %s #config = #iree_codegen.lowering_config<tile_sizes = [[0, 4, 4, 16], [0, 4, 1, 4], [0, 0, 0, 0, 1, 1, 4]]> #translation = #iree_codegen.translation_info<SPIRVVectorize, workload_per_wg = [16, 4, 4]>
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir index 816824b..3d562e2 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-set-num-workgroups,builtin.module(builtin.func(iree-spirv-tile,iree-spirv-vectorize))))' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-set-num-workgroups,builtin.module(func.func(iree-spirv-tile,iree-spirv-vectorize))))' %s | FileCheck %s #config = #iree_codegen.lowering_config<tile_sizes = [[8, 64], [8, 4], [0, 0, 4]]> #translation = #iree_codegen.translation_info<SPIRVVectorize, workload_per_wg = [64, 8]>
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_to_cooperative_ops.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_to_cooperative_ops.mlir index e770765..287ad33 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_to_cooperative_ops.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_to_cooperative_ops.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-tile-and-vectorize-to-cooperative-ops))))' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(func.func(iree-spirv-tile-and-vectorize-to-cooperative-ops))))' %s | FileCheck %s #config = #iree_codegen.lowering_config<tile_sizes = [[16, 16, 16], [16, 16, 16]]> #translation = #iree_codegen.translation_info<SPIRVVectorizeToCooperativeOps, workload_per_wg = [16, 16]>
diff --git a/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir b/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir index 0700af7..301fea3 100644 --- a/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir +++ b/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-vector-to-cooperative-ops,cse))))' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(func.func(iree-spirv-vector-to-cooperative-ops,cse))))' %s | FileCheck %s #map0 = affine_map<(d0, d1) -> (d0, d1)> #map1 = affine_map<(d0, d1, d2) -> (d0, d2)>
diff --git a/iree/compiler/Codegen/Sandbox/test/fusion_expert.mlir b/iree/compiler/Codegen/Sandbox/test/fusion_expert.mlir index beb6930..3ad00e8 100644 --- a/iree/compiler/Codegen/Sandbox/test/fusion_expert.mlir +++ b/iree/compiler/Codegen/Sandbox/test/fusion_expert.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -pass-pipeline="builtin.func(linalg-fuse{tiling-level=0 vectorize}), canonicalize, cse" -split-input-file %s | FileCheck %s +// RUN: iree-opt -pass-pipeline="func.func(linalg-fuse{tiling-level=0 vectorize}), canonicalize, cse" -split-input-file %s | FileCheck %s func @matmul_bias_add(%arg0 : tensor<?x?xf32>, %arg1 : tensor<?x?xf32>, %arg2 : tensor<?xf32>) -> tensor<?x?xf32> { %cst = arith.constant 0.0 : f32
diff --git a/iree/compiler/Codegen/Sandbox/test/single_tiling_expert.mlir b/iree/compiler/Codegen/Sandbox/test/single_tiling_expert.mlir index db90362..9a54ab8 100644 --- a/iree/compiler/Codegen/Sandbox/test/single_tiling_expert.mlir +++ b/iree/compiler/Codegen/Sandbox/test/single_tiling_expert.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -pass-pipeline="builtin.func(linalg-single-tiling-expert-driver{tiling-level=0 vectorize})" -split-input-file %s | FileCheck %s +// RUN: iree-opt -pass-pipeline="func.func(linalg-single-tiling-expert-driver{tiling-level=0 vectorize})" -split-input-file %s | FileCheck %s func @matmul(%arg0 : tensor<?x?xf32>, %arg1 : tensor<?x?xf32>, %arg2 : tensor<?x?xf32>) -> tensor<?x?xf32> { %0 = linalg.matmul {lowering_config = #iree_codegen.lowering_config<tile_sizes = [[10, 20, 30]]>}
diff --git a/iree/compiler/Codegen/Sandbox/test/unroll_one_vector_op.mlir b/iree/compiler/Codegen/Sandbox/test/unroll_one_vector_op.mlir index 43825d1..e015c22 100644 --- a/iree/compiler/Codegen/Sandbox/test/unroll_one_vector_op.mlir +++ b/iree/compiler/Codegen/Sandbox/test/unroll_one_vector_op.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt %s -pass-pipeline='builtin.func(unroll-one-vector-op{anchor-func=test anchor-op=vector.contract source-shape=4,4,3 target-shape=2,4,3})' | FileCheck %s +// RUN: iree-opt %s -pass-pipeline='func.func(unroll-one-vector-op{anchor-func=test anchor-op=vector.contract source-shape=4,4,3 target-shape=2,4,3})' | FileCheck %s #matmul_accesses = [ affine_map<(i, j, k) -> (i, k)>,
diff --git a/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp b/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp index e13b25a..7ce8665 100644 --- a/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp +++ b/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp
@@ -12,6 +12,7 @@ #include "iree/compiler/Codegen/Transforms/Transforms.h" #include "mlir/Dialect/Affine/Utils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/PatternMatch.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" @@ -167,7 +168,7 @@ /// individually. struct AffineMinDistributedSCFCanonicalizationPass : public PassWrapper<AffineMinDistributedSCFCanonicalizationPass, - OperationPass<FuncOp>> { + OperationPass<func::FuncOp>> { StringRef getArgument() const override { return "iree-codegen-affinemin-scf-canonicalization"; } @@ -178,7 +179,7 @@ } void runOnOperation() override { - FuncOp funcOp = getOperation(); + func::FuncOp funcOp = getOperation(); RewritePatternSet foldPattern(&getContext()); populateAffineMinSCFCanonicalizationPattern(foldPattern); FrozenRewritePatternSet frozenPatterns(std::move(foldPattern));
diff --git a/iree/compiler/Codegen/Transforms/Transforms.cpp b/iree/compiler/Codegen/Transforms/Transforms.cpp index 15b258b..a29f1c1 100644 --- a/iree/compiler/Codegen/Transforms/Transforms.cpp +++ b/iree/compiler/Codegen/Transforms/Transforms.cpp
@@ -30,7 +30,7 @@ } // namespace LogicalResult defineWorkgroupCountRegion( - OpBuilder &builder, FuncOp funcOp, + OpBuilder &builder, func::FuncOp funcOp, WorkgroupCountRegionBuilder regionBuilder) { IREE::HAL::ExecutableEntryPointOp entryPointOp = getEntryPoint(funcOp); if (!entryPointOp) {
diff --git a/iree/compiler/Codegen/Transforms/Transforms.h b/iree/compiler/Codegen/Transforms/Transforms.h index cf3fdd2..b25590b 100644 --- a/iree/compiler/Codegen/Transforms/Transforms.h +++ b/iree/compiler/Codegen/Transforms/Transforms.h
@@ -30,7 +30,7 @@ using WorkgroupCountRegionBuilder = std::function<std::array<Value, 3>( OpBuilder &b, Location loc, std::array<Value, 3> workload)>; LogicalResult defineWorkgroupCountRegion( - OpBuilder &builder, FuncOp funcOp, + OpBuilder &builder, func::FuncOp funcOp, WorkgroupCountRegionBuilder regionBuilder); /// Insert patterns to perform folding of AffineMinOp by matching the pattern
diff --git a/iree/compiler/Codegen/Utils/Utils.cpp b/iree/compiler/Codegen/Utils/Utils.cpp index 4651efc..98ad38b 100644 --- a/iree/compiler/Codegen/Utils/Utils.cpp +++ b/iree/compiler/Codegen/Utils/Utils.cpp
@@ -29,9 +29,9 @@ // Utility functions to get entry points //===----------------------------------------------------------------------===// -bool isEntryPoint(FuncOp func) { return func.isPublic(); } +bool isEntryPoint(func::FuncOp func) { return func.isPublic(); } -IREE::HAL::ExecutableEntryPointOp getEntryPoint(FuncOp funcOp) { +IREE::HAL::ExecutableEntryPointOp getEntryPoint(func::FuncOp funcOp) { auto variantOp = funcOp->getParentOfType<IREE::HAL::ExecutableVariantOp>(); for (auto op : variantOp.getOps<IREE::HAL::ExecutableEntryPointOp>()) { if (op.sym_name() == funcOp.getName()) { @@ -436,10 +436,10 @@ } LogicalResult getFilteredOps( - FuncOp funcOp, RootOpFilteringFn filteringFn, + func::FuncOp funcOp, RootOpFilteringFn filteringFn, SmallVectorImpl<Operation *> &filteredOps, SmallVectorImpl<LoopTilingAndDistributionInfo> &tiledLoops) { - Region ®ion = funcOp.body(); + Region ®ion = funcOp.getBody(); if (!llvm::hasSingleElement(region)) { return funcOp.emitError("unable dispatch function with multiple blocks"); } @@ -463,7 +463,7 @@ } LogicalResult getComputeOps( - FuncOp funcOp, SmallVectorImpl<Operation *> &computeOps, + func::FuncOp funcOp, SmallVectorImpl<Operation *> &computeOps, SmallVectorImpl<LoopTilingAndDistributionInfo> &tiledLoops) { if (failed(getFilteredOps( funcOp, @@ -477,7 +477,7 @@ } SmallVector<LoopTilingAndDistributionInfo> getTiledAndDistributedLoopInfo( - FuncOp funcOp) { + func::FuncOp funcOp) { SmallVector<LoopTilingAndDistributionInfo> info; funcOp.walk([&](scf::ForOp forOp) { if (auto tiledLoopInfo = isTiledAndDistributedLoop(forOp)) {
diff --git a/iree/compiler/Codegen/Utils/Utils.h b/iree/compiler/Codegen/Utils/Utils.h index 5110269..e137eb1 100644 --- a/iree/compiler/Codegen/Utils/Utils.h +++ b/iree/compiler/Codegen/Utils/Utils.h
@@ -11,6 +11,7 @@ #include "iree/compiler/Dialect/HAL/IR/HALOps.h" #include "llvm/ADT/StringMap.h" #include "llvm/ADT/Triple.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/SCF/SCF.h" #include "mlir/IR/BuiltinOps.h" @@ -25,24 +26,24 @@ //===----------------------------------------------------------------------===// /// Returns true if the given `func` is a kernel dispatch entry point. -bool isEntryPoint(FuncOp func); +bool isEntryPoint(func::FuncOp func); /// Returns a map from function symbol name to corresponding entry point op. llvm::StringMap<IREE::HAL::ExecutableEntryPointOp> getAllEntryPoints( ModuleOp module); /// Returns the entry point op for the `funcOp`. Returns `nullptr` on failure. -IREE::HAL::ExecutableEntryPointOp getEntryPoint(FuncOp funcOp); +IREE::HAL::ExecutableEntryPointOp getEntryPoint(func::FuncOp funcOp); /// Methods to get backend information. bool isX86(IREE::HAL::ExecutableVariantOp variantOp); -inline bool isX86(FuncOp entryPointFn) { +inline bool isX86(func::FuncOp entryPointFn) { auto variantOp = entryPointFn->getParentOfType<IREE::HAL::ExecutableVariantOp>(); return isX86(variantOp); } bool isRISCV(IREE::HAL::ExecutableVariantOp variantOp); -inline bool isRISCV(FuncOp entryPointFn) { +inline bool isRISCV(func::FuncOp entryPointFn) { auto variantOp = entryPointFn->getParentOfType<IREE::HAL::ExecutableVariantOp>(); return isRISCV(variantOp); @@ -50,7 +51,7 @@ inline bool isVMVXBackend(IREE::HAL::ExecutableVariantOp variantOp) { return variantOp.target().getBackend().getValue() == "vmvx"; } -inline bool isVMVXBackend(FuncOp entryPointFn) { +inline bool isVMVXBackend(func::FuncOp entryPointFn) { auto variantOp = entryPointFn->getParentOfType<IREE::HAL::ExecutableVariantOp>(); return isVMVXBackend(variantOp); @@ -122,14 +123,14 @@ /// other cases. using RootOpFilteringFn = std::function<bool(Operation *)>; LogicalResult getFilteredOps( - FuncOp funcOp, RootOpFilteringFn filteringFn, + func::FuncOp funcOp, RootOpFilteringFn filteringFn, SmallVectorImpl<Operation *> &filteredOps, SmallVectorImpl<LoopTilingAndDistributionInfo> &tiledLoops); /// Specialization of `getFilteredOps` for filtering `LinalgOp`s and /// `LinagExtOp`s. LogicalResult getComputeOps( - FuncOp funcOp, SmallVectorImpl<Operation *> &computeOps, + func::FuncOp funcOp, SmallVectorImpl<Operation *> &computeOps, SmallVectorImpl<LoopTilingAndDistributionInfo> &tiledLoops); /// If the given `forOp` is a tiled and distributed loop, returns its tiling and @@ -139,7 +140,7 @@ /// Collects information about loops matching tiled+distribute pattern. SmallVector<LoopTilingAndDistributionInfo> getTiledAndDistributedLoopInfo( - FuncOp funcOp); + func::FuncOp funcOp); Operation *createLinalgCopyOp(OpBuilder &b, Location loc, Value from, Value to, ArrayRef<NamedAttribute> attributes = {});
diff --git a/iree/compiler/Dialect/Flow/IR/PartitionableLoopsInterface.cpp b/iree/compiler/Dialect/Flow/IR/PartitionableLoopsInterface.cpp index b6aca9e..1f245bf 100644 --- a/iree/compiler/Dialect/Flow/IR/PartitionableLoopsInterface.cpp +++ b/iree/compiler/Dialect/Flow/IR/PartitionableLoopsInterface.cpp
@@ -6,6 +6,7 @@ #include "iree/compiler/Dialect/Flow/IR/PartitionableLoopsInterface.h" +#include "iree-dialects/Dialect/LinalgExt/IR/LinalgExtDialect.h" #include "iree-dialects/Dialect/LinalgExt/IR/LinalgExtOps.h" #include "iree-dialects/Dialect/LinalgExt/IR/TiledOpInterface.h" #include "llvm/ADT/SmallVector.h" @@ -131,40 +132,43 @@ /// needs to be done on a op-by-op basis since registration is on an op-by-op /// basis. template <typename OpTy> -static void registerInterfaceForLinalgOps(DialectRegistry ®istry) { - registry.addOpInterface<OpTy, LinalgOpPartitionableLoops>(); +static void registerInterfaceForLinalgOps(MLIRContext *ctx) { + OpTy::template attachInterface<LinalgOpPartitionableLoops>(*ctx); } /// Specializations of the registration method to use a different external model /// instead of the generic external model for Linalg ops. template <> -void registerInterfaceForLinalgOps<linalg::Mmt4DOp>(DialectRegistry ®istry) { - registry.addOpInterface<linalg::Mmt4DOp, Mmt4DOpPartitionableLoops>(); +void registerInterfaceForLinalgOps<linalg::Mmt4DOp>(MLIRContext *ctx) { + linalg::Mmt4DOp::attachInterface<Mmt4DOpPartitionableLoops>(*ctx); } /// Registers the external models for all Linalg operations. template <typename OpTy1, typename OpTy2, typename... More> -static void registerInterfaceForLinalgOps(DialectRegistry ®istry) { - registerInterfaceForLinalgOps<OpTy1>(registry); - registerInterfaceForLinalgOps<OpTy2, More...>(registry); +static void registerInterfaceForLinalgOps(MLIRContext *ctx) { + registerInterfaceForLinalgOps<OpTy1>(ctx); + registerInterfaceForLinalgOps<OpTy2, More...>(ctx); } /// Registers the `TiledOpInterfacePartitionableLoops` model for operations. template <typename OpTy> -static void registerInterfaceForTiledOpInterfaceOps(DialectRegistry ®istry) { - registry.addOpInterface<OpTy, TiledOpInterfacePartitionableLoops>(); +static void registerInterfaceForTiledOpInterfaceOps(MLIRContext *ctx) { + OpTy ::template attachInterface<TiledOpInterfacePartitionableLoops>(*ctx); } /// Registers the external models for all TiledOpInterface operations. template <typename OpTy1, typename OpTy2, typename... More> -static void registerInterfaceForTiledOpInterfaceOps(DialectRegistry ®istry) { - registerInterfaceForTiledOpInterfaceOps<OpTy1>(registry); - registerInterfaceForTiledOpInterfaceOps<OpTy2, More...>(registry); +static void registerInterfaceForTiledOpInterfaceOps(MLIRContext *ctx) { + registerInterfaceForTiledOpInterfaceOps<OpTy1>(ctx); + registerInterfaceForTiledOpInterfaceOps<OpTy2, More...>(ctx); } void registerPartitionableLoopsInterfaceModels(DialectRegistry ®istry) { - // clang-format off - registerInterfaceForLinalgOps< + registry.insert<linalg::LinalgDialect>(); + + registry.addExtension(+[](MLIRContext *ctx, linalg::LinalgDialect *dialect) { + registerInterfaceForLinalgOps< + // clang-format off // This is copy-pasted from LinalgStructuredOps.cpp.inc. In theory you could // just include that generated file here, but that cause errors with bazel. @@ -206,13 +210,22 @@ ::mlir::linalg::QuantizedBatchMatmulOp, ::mlir::linalg::QuantizedMatmulOp, ::mlir::linalg::VecmatOp - >(registry); - // clang-format on + // clang-format on + >(ctx); + }); - registerInterfaceForTiledOpInterfaceOps< - LinalgExt::FftOp, LinalgExt::ReverseOp, LinalgExt::ScanOp, - LinalgExt::ScatterOp, LinalgExt::SortOp, tensor::ExtractSliceOp, - tensor::InsertSliceOp>(registry); + registry.insert<LinalgExt::IREELinalgExtDialect>(); + + registry.addExtension( + +[](MLIRContext *ctx, LinalgExt::IREELinalgExtDialect *dialect) { + registerInterfaceForTiledOpInterfaceOps< + LinalgExt::FftOp, LinalgExt::ReverseOp, LinalgExt::ScanOp, + LinalgExt::ScatterOp, LinalgExt::SortOp>(ctx); + }); + registry.addExtension(+[](MLIRContext *ctx, tensor::TensorDialect *dialect) { + registerInterfaceForTiledOpInterfaceOps<tensor::ExtractSliceOp, + tensor::InsertSliceOp>(ctx); + }); } } // namespace Flow
diff --git a/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp b/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp index 0e77a72..738b1b2 100644 --- a/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp
@@ -151,8 +151,8 @@ }; } // namespace -std::unique_ptr<OperationPass<mlir::FuncOp>> createConvertLinalgTensorOpsPass( - bool runBeforeDispatchRegionFormation) { +std::unique_ptr<OperationPass<mlir::func::FuncOp>> +createConvertLinalgTensorOpsPass(bool runBeforeDispatchRegionFormation) { return std::make_unique<ConvertLinalgTensorOpsPass>( runBeforeDispatchRegionFormation); }
diff --git a/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp b/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp index b33cdab..902b65a 100644 --- a/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp
@@ -207,13 +207,15 @@ } // Walk all functions and ensure equivalent. - if (!compare_ranges( - lhsModule.getOps<mlir::FuncOp>(), rhsModule.getOps<mlir::FuncOp>(), - [](mlir::FuncOp lhs, mlir::FuncOp rhs) { - if (lhs.getType() != rhs.getType()) return false; - if (lhs->getAttrs() != rhs->getAttrs()) return false; - return isStructurallyEquivalentTo(lhs.getRegion(), rhs.getRegion()); - })) { + if (!compare_ranges(lhsModule.getOps<mlir::func::FuncOp>(), + rhsModule.getOps<mlir::func::FuncOp>(), + [](mlir::func::FuncOp lhs, mlir::func::FuncOp rhs) { + if (lhs.getFunctionType() != rhs.getFunctionType()) + return false; + if (lhs->getAttrs() != rhs->getAttrs()) return false; + return isStructurallyEquivalentTo(lhs.getRegion(), + rhs.getRegion()); + })) { return false; // dispatch entry mismatch }
diff --git a/iree/compiler/Dialect/Flow/Transforms/ExpandTensorShapes.cpp b/iree/compiler/Dialect/Flow/Transforms/ExpandTensorShapes.cpp index a560f70..2c08f14 100644 --- a/iree/compiler/Dialect/Flow/Transforms/ExpandTensorShapes.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/ExpandTensorShapes.cpp
@@ -321,9 +321,9 @@ // -> // func @foo(%0: tensor<?xf32>, %d: index) { // %1 = flow.tensor.tie_shape %0 : tensor<?xf32>{%d} -static void expandFuncOp(mlir::FuncOp op, ExpandedGlobalMap &globalMap, +static void expandFuncOp(mlir::func::FuncOp op, ExpandedGlobalMap &globalMap, IndexSet &indexSet, TensorDimMap &tensorDimMap) { - auto oldType = op.getType(); + auto oldType = op.getFunctionType(); auto inputTypes = expandTypes(oldType.getInputs()); auto resultTypes = expandTypes(oldType.getResults()); auto newType = FunctionType::get(op.getContext(), inputTypes, resultTypes); @@ -482,7 +482,7 @@ expandGlobalStoreOp(storeOp, globalMap, indexSet, tensorDimMap); } else if (auto initializerOp = dyn_cast<IREE::Util::InitializerOp>(op)) { expandInitializerOp(initializerOp, globalMap, indexSet, tensorDimMap); - } else if (auto funcOp = dyn_cast<mlir::FuncOp>(op)) { + } else if (auto funcOp = dyn_cast<mlir::func::FuncOp>(op)) { expandFuncOp(funcOp, globalMap, indexSet, tensorDimMap); } else if (auto callOp = dyn_cast<mlir::func::CallOp>(op)) { expandCallOp(callOp, indexSet, tensorDimMap);
diff --git a/iree/compiler/Dialect/Flow/Transforms/ExportBenchmarkFuncs.cpp b/iree/compiler/Dialect/Flow/Transforms/ExportBenchmarkFuncs.cpp index 0fa5bcd..61504aa 100644 --- a/iree/compiler/Dialect/Flow/Transforms/ExportBenchmarkFuncs.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/ExportBenchmarkFuncs.cpp
@@ -37,8 +37,8 @@ // Gather the functions we want to wrap for benchmarking and wrap them. // Since we are inserting new functions as part of this pass we must perform // the wrapping for only the inputs. - SmallVector<mlir::FuncOp, 4> entryFuncOps; - for (auto entryFuncOp : moduleOp.getOps<mlir::FuncOp>()) { + SmallVector<mlir::func::FuncOp, 4> entryFuncOps; + for (auto entryFuncOp : moduleOp.getOps<mlir::func::FuncOp>()) { if (entryFuncOp.isPublic()) { entryFuncOps.push_back(entryFuncOp); } @@ -71,14 +71,14 @@ } LogicalResult createEntryPointBenchmarkFunc(mlir::ModuleOp moduleOp, - mlir::FuncOp entryFuncOp) { + mlir::func::FuncOp entryFuncOp) { OpBuilder moduleBuilder(&getContext()); moduleBuilder.setInsertionPointAfter(entryFuncOp); // Create one dummy input variable per input. Location loc = entryFuncOp.getLoc(); SmallVector<IREE::Util::GlobalOp, 4> dummyInputVariableOps; - for (auto inputType : entryFuncOp.getType().getInputs()) { + for (auto inputType : entryFuncOp.getFunctionType().getInputs()) { auto dummyVar = createDummyInputVariableOp(loc, inputType, moduleBuilder); if (!dummyVar) return failure(); dummyInputVariableOps.push_back(dummyVar); @@ -86,7 +86,7 @@ // Create a `() -> ()` entry point op the benchmark tool can run. std::string funcName = std::string(entryFuncOp.getName()) + "_benchmark"; - auto funcOp = moduleBuilder.create<mlir::FuncOp>( + auto funcOp = moduleBuilder.create<mlir::func::FuncOp>( loc, funcName, moduleBuilder.getFunctionType({}, {})); funcOp.setPublic(); funcOp->setAttr("iree.abi.stub", moduleBuilder.getUnitAttr());
diff --git a/iree/compiler/Dialect/Flow/Transforms/OutlineDispatchRegions.cpp b/iree/compiler/Dialect/Flow/Transforms/OutlineDispatchRegions.cpp index 43d51f2..2a7e7af 100644 --- a/iree/compiler/Dialect/Flow/Transforms/OutlineDispatchRegions.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/OutlineDispatchRegions.cpp
@@ -30,7 +30,7 @@ // Creates a flow.executable out of a set of functions, pulling in all other // functions reachable by the provided functions. static ExecutableOp createExecutable(Location loc, StringRef executableName, - ArrayRef<mlir::FuncOp> funcOps, + ArrayRef<mlir::func::FuncOp> funcOps, ModuleOp parentModuleOp) { assert(!funcOps.empty() && "must have at least one entry function"); @@ -86,14 +86,15 @@ } // Converts a dispatch region body to a free-floating function. -static mlir::FuncOp createWorkgroupFunc(Location loc, StringRef functionName, - Region ®ion) { +static mlir::func::FuncOp createWorkgroupFunc(Location loc, + StringRef functionName, + Region ®ion) { // Build function type matching the region signature. auto functionType = FunctionType::get( region.getContext(), region.getArgumentTypes(), /*results=*/{}); // Clone region into the function body. - auto funcOp = mlir::FuncOp::create(loc, functionName, functionType); + auto funcOp = mlir::func::FuncOp::create(loc, functionName, functionType); BlockAndValueMapping mapping; region.cloneInto(&funcOp.getBody(), mapping); @@ -158,7 +159,7 @@ // Generate a nice name if possible. std::string opName; - if (auto funcOp = llvm::dyn_cast<mlir::FuncOp>(operation)) { + if (auto funcOp = llvm::dyn_cast<mlir::func::FuncOp>(operation)) { opName = funcOp.getName().str(); } else if (llvm::isa<IREE::Util::InitializerOp>(operation)) { opName =
diff --git a/iree/compiler/Dialect/Flow/Transforms/PassDetail.h b/iree/compiler/Dialect/Flow/Transforms/PassDetail.h index 0c80906..740ad78 100644 --- a/iree/compiler/Dialect/Flow/Transforms/PassDetail.h +++ b/iree/compiler/Dialect/Flow/Transforms/PassDetail.h
@@ -7,6 +7,7 @@ #ifndef IREE_COMPILER_DIALECT_FLOW_TRANSFORMS_PASS_DETAIL_H_ #define IREE_COMPILER_DIALECT_FLOW_TRANSFORMS_PASS_DETAIL_H_ +#include "mlir/IR/FunctionInterfaces.h" #include "mlir/Pass/Pass.h" namespace mlir {
diff --git a/iree/compiler/Dialect/Flow/Transforms/Passes.cpp b/iree/compiler/Dialect/Flow/Transforms/Passes.cpp index d6417bc..7e08332 100644 --- a/iree/compiler/Dialect/Flow/Transforms/Passes.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/Passes.cpp
@@ -10,6 +10,7 @@ #include "iree/compiler/Dialect/Util/Transforms/Passes.h" #include "iree/compiler/Utils/PassUtils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Dialect/MemRef/Transforms/Passes.h" #include "mlir/Pass/PassOptions.h" @@ -79,7 +80,7 @@ namespace { -using FunctionLikeNest = MultiOpNest<FuncOp, IREE::Util::InitializerOp>; +using FunctionLikeNest = MultiOpNest<func::FuncOp, IREE::Util::InitializerOp>; // Subset of the overall pass pipeline for optimizing globals and numerics. // We may ultimately break this out separately so creating a syntactic
diff --git a/iree/compiler/Dialect/Flow/Transforms/PromoteTensorLoads.cpp b/iree/compiler/Dialect/Flow/Transforms/PromoteTensorLoads.cpp index c67ab91..afa2a64 100644 --- a/iree/compiler/Dialect/Flow/Transforms/PromoteTensorLoads.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/PromoteTensorLoads.cpp
@@ -91,7 +91,8 @@ } }; -std::unique_ptr<OperationPass<mlir::FuncOp>> createPromoteTensorLoadsPass() { +std::unique_ptr<OperationPass<mlir::func::FuncOp>> +createPromoteTensorLoadsPass() { return std::make_unique<PromoteTensorLoadsPass>(); }
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir b/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir index 6d546da..b83c34c 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir
@@ -162,7 +162,7 @@ flow.dispatch.entry @different_types_float_entry builtin.module { func @different_types_float_entry(%arg0: tensor<4xf32>) -> tensor<4xi1> { - %0 = "mhlo.compare"(%arg0, %arg0) {comparison_direction = "EQ"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xi1> + %0 = "mhlo.compare"(%arg0, %arg0) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xi1> return %0 : tensor<4xi1> } } @@ -172,7 +172,7 @@ flow.dispatch.entry @different_types_int_entry builtin.module { func @different_types_int_entry(%arg0: tensor<4xi32>) -> tensor<4xi1> { - %0 = "mhlo.compare"(%arg0, %arg0) {comparison_direction = "EQ"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %0 = "mhlo.compare"(%arg0, %arg0) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> return %0 : tensor<4xi1> } }
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir index d74946a..59553df 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -verify-diagnostics -pass-pipeline="builtin.func(iree-flow-dispatch-linalg-on-tensors-pass, resolve-shaped-type-result-dims, cse, canonicalize, cse)" %s | FileCheck %s +// RUN: iree-opt -split-input-file -verify-diagnostics -pass-pipeline="func.func(iree-flow-dispatch-linalg-on-tensors-pass, resolve-shaped-type-result-dims, cse, canonicalize, cse)" %s | FileCheck %s func @tile_matmul_alone(%arg0 : tensor<?x?xf32>, %arg1 : tensor<?x?xf32>, %arg2 : tensor<?x?xf32>) -> tensor<?x?xf32> {
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_fusion.mlir b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_fusion.mlir index d24826e..74fc741 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_fusion.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_fusion.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -verify-diagnostics -pass-pipeline="builtin.func(iree-flow-dispatch-linalg-on-tensors-pass)" -canonicalize -cse %s | FileCheck %s +// RUN: iree-opt -split-input-file -verify-diagnostics -pass-pipeline="func.func(iree-flow-dispatch-linalg-on-tensors-pass)" -canonicalize -cse %s | FileCheck %s func @fuse_conv2d_elementwise(%input: tensor<1x225x225x16xf32>, %filter: tensor<3x3x16x32xf32>, %offset: tensor<32xf32>) -> tensor<1x112x112x32xf32> { %cst = arith.constant 0.000000e+00 : f32
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/export_benchmark_funcs.mlir b/iree/compiler/Dialect/Flow/Transforms/test/export_benchmark_funcs.mlir index dbc5ff7..9bf29cc 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/export_benchmark_funcs.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/export_benchmark_funcs.mlir
@@ -22,7 +22,7 @@ func @while(%start: tensor<i32>, %bound: tensor<i32>) -> tensor<i32> { cf.br ^bb1(%start : tensor<i32>) ^bb1(%0: tensor<i32>): - %1 = "mhlo.compare"(%0, %bound) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %1 = "mhlo.compare"(%0, %bound) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> %2 = tensor.extract %1[] : tensor<i1> cf.cond_br %2, ^bb2(%0 : tensor<i32>), ^bb3(%0 : tensor<i32>) ^bb2(%3: tensor<i32>):
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/inject_dispatch_tracing.mlir b/iree/compiler/Dialect/Flow/Transforms/test/inject_dispatch_tracing.mlir index b029c54..04500aa 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/inject_dispatch_tracing.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/inject_dispatch_tracing.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='builtin.func(iree-flow-inject-dispatch-tracing)' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='func.func(iree-flow-inject-dispatch-tracing)' %s | FileCheck %s // CHECK-LABEL: func @singleDispatch // CHECK-SAME: (%[[ARG0:.+]]: tensor<4xf32>)
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/strip_signedness.mlir b/iree/compiler/Dialect/Flow/Transforms/test/strip_signedness.mlir index 80a1bf1..7bd6c3b 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/strip_signedness.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/strip_signedness.mlir
@@ -1,5 +1,5 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='builtin.func(iree-flow-strip-signedness)' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='func.func(iree-flow-strip-signedness)' %s | FileCheck %s // CHECK-LABEL: @strip_signedness_arg // CHECK-SAME: tensor<4xi8>
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/verify_input_ir.mlir b/iree/compiler/Dialect/Flow/Transforms/test/verify_input_ir.mlir index 6f678bd..1400b00 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/verify_input_ir.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/verify_input_ir.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -pass-pipeline="builtin.func(iree-verify-input-legality)" -verify-diagnostics %s -split-input-file +// RUN: iree-opt -pass-pipeline="func.func(iree-verify-input-legality)" -verify-diagnostics %s -split-input-file // expected-error@below {{illegal operations still remain}} func @check_no_mhlo(%arg0: tensor<?x?xf32>, %arg1 : tensor<?x?xf32>) -> tensor<?x?xf32> {
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferOps.cpp b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferOps.cpp index 1b7a323..99112ae 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferOps.cpp
@@ -25,7 +25,7 @@ LogicalResult matchAndRewrite( IREE::HAL::BufferLoadOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - auto importType = importOp.getType(); + auto importType = importOp.getFunctionType(); auto originalType = op.result().getType(); auto targetType = typeConverter->convertType(op.result().getType()); @@ -105,7 +105,7 @@ LogicalResult matchAndRewrite( IREE::HAL::BufferStoreOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - auto importType = importOp.getType(); + auto importType = importOp.getFunctionType(); auto elementType = op.value().getType(); int32_t validByteWidth =
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp index 6ded056..7807747 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp
@@ -29,7 +29,7 @@ LogicalResult matchAndRewrite( IREE::HAL::CommandBufferFillBufferOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - auto importType = importOp.getType(); + auto importType = importOp.getFunctionType(); SmallVector<Value, 8> callOperands = { adaptor.command_buffer(), @@ -84,7 +84,7 @@ LogicalResult matchAndRewrite( IREE::HAL::CommandBufferPushDescriptorSetOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - auto importType = importOp.getType(); + auto importType = importOp.getFunctionType(); SmallVector<Value, 8> callOperands = { adaptor.command_buffer(),
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertExecutableOps.cpp b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertExecutableOps.cpp index 8752f9f..5094543 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertExecutableOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertExecutableOps.cpp
@@ -77,7 +77,7 @@ auto executableFormatString = detail::rewriteAttrToOperands( createOp.getLoc(), executableBinaryOp.formatAttr(), - importOp.getType().getInput(1), rewriter); + importOp.getFunctionType().getInput(1), rewriter); assert(executableFormatString.hasValue() && executableFormatString.getValue().size() == 1); auto executableRodata = @@ -103,7 +103,7 @@ }; callOperands.append(adaptor.layouts().begin(), adaptor.layouts().end()); - auto importType = importOp.getType(); + auto importType = importOp.getFunctionType(); auto callOp = rewriter.replaceOpWithNewOp<IREE::VM::CallVariadicOp>( createOp, SymbolRefAttr::get(importOp), importType.getResults(), segmentSizes, importType.getInputs(), callOperands);
diff --git a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStructuralOps.cpp b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStructuralOps.cpp index e18b1d0..5e1077f 100644 --- a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStructuralOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStructuralOps.cpp
@@ -23,18 +23,19 @@ namespace iree_compiler { namespace { -class FuncOpSignatureConversion : public OpConversionPattern<mlir::FuncOp> { +class FuncOpSignatureConversion + : public OpConversionPattern<mlir::func::FuncOp> { public: using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - mlir::FuncOp funcOp, OpAdaptor adaptor, + mlir::func::FuncOp funcOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { auto &typeConverter = *getTypeConverter(); // Convert the input signature types. // TODO(benvanik): dynamic shapes by passing in tensor dynamic dims. - auto originalType = funcOp.getType(); + auto originalType = funcOp.getFunctionType(); TypeConverter::SignatureConversion newSignature( originalType.getNumInputs()); for (auto argType : llvm::enumerate(originalType.getInputs())) { @@ -185,10 +186,11 @@ // We need to rewrite certain types on operands/results so use the default // dynamic legality checker to force any ops using such types to run through // our patterns. - conversionTarget.addDynamicallyLegalOp<mlir::FuncOp>([&](mlir::FuncOp op) { - return typeConverter.isSignatureLegal(op.getType()) && - typeConverter.isLegal(&op.getBody()); - }); + conversionTarget.addDynamicallyLegalOp<mlir::func::FuncOp>( + [&](mlir::func::FuncOp op) { + return typeConverter.isSignatureLegal(op.getFunctionType()) && + typeConverter.isLegal(&op.getBody()); + }); patterns .insert<FuncOpSignatureConversion, CallOpConversion, ReturnOpConversion,
diff --git a/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp b/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp index 7d99766..d782458 100644 --- a/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp +++ b/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp
@@ -98,8 +98,9 @@ //===----------------------------------------------------------------------===// // Verifies that all types used with the given entry point are supportable. -static LogicalResult verifyEntryPointTypes(mlir::FuncOp entryFuncOp) { - for (auto inputType : llvm::enumerate(entryFuncOp.getType().getInputs())) { +static LogicalResult verifyEntryPointTypes(mlir::func::FuncOp entryFuncOp) { + for (auto inputType : + llvm::enumerate(entryFuncOp.getFunctionType().getInputs())) { if (inputType.value().isa<IREE::Stream::BindingType>() || inputType.value().isInteger(32)) { // OK - directly translates to a HAL interface binding. @@ -131,8 +132,9 @@ } // Converts the usage of the given primitive |arg| to interface methods. -static void convertOperandUsage(mlir::FuncOp sourceFuncOp, BlockArgument arg, - unsigned pushConstantIdx, OpBuilder &builder) { +static void convertOperandUsage(mlir::func::FuncOp sourceFuncOp, + BlockArgument arg, unsigned pushConstantIdx, + OpBuilder &builder) { auto alignmentAttr = sourceFuncOp.getArgAttrOfType<IntegerAttr>( arg.getArgNumber(), "stream.alignment"); auto valuesAttr = sourceFuncOp.getArgAttrOfType<ArrayAttr>(arg.getArgNumber(), @@ -145,7 +147,7 @@ // Converts the usage of the given !stream.binding |arg| to interface methods. static void convertBindingUsage( - mlir::FuncOp sourceFuncOp, BlockArgument arg, + mlir::func::FuncOp sourceFuncOp, BlockArgument arg, IREE::HAL::DescriptorSetLayoutAttr setLayoutAttr, IREE::HAL::DescriptorSetBindingAttr bindingAttr) { if (arg.use_empty()) return; // no-op @@ -166,8 +168,8 @@ // Clones |sourceFuncOp| and updates its signature to match the |interfaceOp| // and use the HAL interface access primitives. -static mlir::FuncOp cloneFuncWithInterface( - mlir::FuncOp sourceFuncOp, const ExecutableLayout &executableLayout, +static mlir::func::FuncOp cloneFuncWithInterface( + mlir::func::FuncOp sourceFuncOp, const ExecutableLayout &executableLayout, IREE::HAL::ExecutableLayoutAttr layoutAttr) { // Clone so that we can do a bunch of unsafe in-place updates. auto clonedFuncOp = sourceFuncOp.clone(); @@ -231,7 +233,7 @@ sourceExecutableOp.body().getOps<IREE::Stream::ExecutableExportOp>()) { int ordinal = nextOrdinal++; auto sourceFuncOp = - sourceExecutableOp.getInnerModule().lookupSymbol<mlir::FuncOp>( + sourceExecutableOp.getInnerModule().lookupSymbol<mlir::func::FuncOp>( exportOp.function_ref()); if (failed(verifyEntryPointTypes(sourceFuncOp))) return failure(); @@ -296,7 +298,7 @@ LogicalResult matchAndRewrite(IREE::HAL::InterfaceWorkgroupSizeOp sizeOp, PatternRewriter &rewriter) const override { // Lookup the entry point matching the parent. - auto funcOp = sizeOp->getParentOfType<mlir::FuncOp>(); + auto funcOp = sizeOp->getParentOfType<mlir::func::FuncOp>(); auto variantOp = funcOp->getParentOfType<IREE::HAL::ExecutableVariantOp>(); auto entryPointOp = dyn_cast<IREE::HAL::ExecutableEntryPointOp>( SymbolTable::lookupSymbolIn(variantOp, funcOp.getName()));
diff --git a/iree/compiler/Dialect/HAL/Transforms/PackDispatchOperands.cpp b/iree/compiler/Dialect/HAL/Transforms/PackDispatchOperands.cpp index dc03c3c..1de050f 100644 --- a/iree/compiler/Dialect/HAL/Transforms/PackDispatchOperands.cpp +++ b/iree/compiler/Dialect/HAL/Transforms/PackDispatchOperands.cpp
@@ -134,7 +134,7 @@ // that was applied to dispatch ops above. // // This is a mirror of updateDispatchOp; see that for more information. -static void updateExportFuncOp(mlir::FuncOp funcOp) { +static void updateExportFuncOp(mlir::func::FuncOp funcOp) { assert(!funcOp.empty() && "can't have empty exported functions"); auto &entryBlock = funcOp.getBody().front(); auto builder = OpBuilder::atBlockBegin(&entryBlock); @@ -225,8 +225,8 @@ } if (newArgTypes.size() != funcOp.getNumArguments()) { // Changed argument count; update signature. - funcOp.setType( - builder.getFunctionType(newArgTypes, funcOp.getType().getResults())); + funcOp.setType(builder.getFunctionType( + newArgTypes, funcOp.getFunctionType().getResults())); funcOp.setAllArgAttrs(newArgAttrs); } @@ -289,8 +289,8 @@ newArgTypes.push_back(builder.getI32Type()); } - funcOp.setType( - builder.getFunctionType(newArgTypes, funcOp.getType().getResults())); + funcOp.setType(builder.getFunctionType( + newArgTypes, funcOp.getFunctionType().getResults())); } //===----------------------------------------------------------------------===// @@ -319,7 +319,8 @@ // Convert all public function signatures and manipulate the arguments. for (auto executableOp : getOperation().getOps<IREE::Stream::ExecutableOp>()) { - for (auto funcOp : executableOp.getInnerModule().getOps<mlir::FuncOp>()) { + for (auto funcOp : + executableOp.getInnerModule().getOps<mlir::func::FuncOp>()) { if (funcOp.isPublic()) { updateExportFuncOp(funcOp); }
diff --git a/iree/compiler/Dialect/HAL/Transforms/Passes.cpp b/iree/compiler/Dialect/HAL/Transforms/Passes.cpp index 13e3edb..34787c7 100644 --- a/iree/compiler/Dialect/HAL/Transforms/Passes.cpp +++ b/iree/compiler/Dialect/HAL/Transforms/Passes.cpp
@@ -59,7 +59,7 @@ // redundant store-loads are removed. passManager.addNestedPass<IREE::Util::InitializerOp>( IREE::Util::createSimplifyGlobalAccessesPass()); - passManager.addNestedPass<mlir::FuncOp>( + passManager.addNestedPass<mlir::func::FuncOp>( IREE::Util::createSimplifyGlobalAccessesPass()); // Cleanup and canonicalization of util.global (and other util ops). @@ -169,7 +169,8 @@ // better CSE/fold dispatch logic. passManager.addNestedPass<IREE::Util::InitializerOp>( createInlineDeviceSwitchesPass()); - passManager.addNestedPass<mlir::FuncOp>(createInlineDeviceSwitchesPass()); + passManager.addNestedPass<mlir::func::FuncOp>( + createInlineDeviceSwitchesPass()); // Memoize device queries such that we don't need to repeatedly ask the same // information at runtime. @@ -180,14 +181,15 @@ // HACK: repeat dispatch ops for benchmarks. if (benchmarkDispatchRepeatCount != 1) { - passManager.addNestedPass<mlir::FuncOp>( + passManager.addNestedPass<mlir::func::FuncOp>( createBenchmarkBatchDispatchesPass(benchmarkDispatchRepeatCount)); } // Elide redundant command buffer state ops created during conversion. passManager.addNestedPass<IREE::Util::InitializerOp>( createElideRedundantCommandsPass()); - passManager.addNestedPass<mlir::FuncOp>(createElideRedundantCommandsPass()); + passManager.addNestedPass<mlir::func::FuncOp>( + createElideRedundantCommandsPass()); // Fixup workgroup count calculations that may have used the affine dialect. // Kind of random here but can happen if the benchmarking code does things.
diff --git a/iree/compiler/Dialect/HAL/Transforms/Passes.h b/iree/compiler/Dialect/HAL/Transforms/Passes.h index 8bcda15..3b47c2b 100644 --- a/iree/compiler/Dialect/HAL/Transforms/Passes.h +++ b/iree/compiler/Dialect/HAL/Transforms/Passes.h
@@ -10,6 +10,7 @@ #include "iree/compiler/Dialect/HAL/IR/HALOps.h" #include "iree/compiler/Dialect/HAL/Target/TargetBackend.h" #include "llvm/ADT/StringMap.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/Pass/Pass.h" #include "mlir/Pass/PassManager.h" @@ -111,7 +112,7 @@ //===----------------------------------------------------------------------===// // Performs packing and materializes runtime packing code when required. -std::unique_ptr<OperationPass<FuncOp>> createPackAllocationsPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createPackAllocationsPass(); // Finds all resource lookups (such as hal.executable.lookup), materializes // their cache storage and initialization, and rewrites the lookups to @@ -121,14 +122,14 @@ // Eliminates redundant 'load's of variables within functions with no 'store'. // TODO(#1124): replace with memory side effects once supported upstream. -std::unique_ptr<OperationPass<FuncOp>> createCSEVariableLoadsPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createCSEVariableLoadsPass(); // Elides stateful command buffer ops that set redundant state. std::unique_ptr<OperationPass<void>> createElideRedundantCommandsPass(); // Repeats dispatches `iree-hal-repeat-dispatch-num` times, which is 1 by // default. -std::unique_ptr<OperationPass<FuncOp>> createBenchmarkBatchDispatchesPass( +std::unique_ptr<OperationPass<func::FuncOp>> createBenchmarkBatchDispatchesPass( unsigned repeatCount); //===----------------------------------------------------------------------===//
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/elide_redundant_commands.mlir b/iree/compiler/Dialect/HAL/Transforms/test/elide_redundant_commands.mlir index 3184ed3..68c08e2 100644 --- a/iree/compiler/Dialect/HAL/Transforms/test/elide_redundant_commands.mlir +++ b/iree/compiler/Dialect/HAL/Transforms/test/elide_redundant_commands.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='builtin.func(iree-hal-elide-redundant-commands)' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='func.func(iree-hal-elide-redundant-commands)' %s | FileCheck %s // Tests that redundant barriers are elided but barriers gaurding ops are not.
diff --git a/iree/compiler/Dialect/HAL/Utils/InferCustomKernelsTargetInfoFromParent.cpp b/iree/compiler/Dialect/HAL/Utils/InferCustomKernelsTargetInfoFromParent.cpp index 15d31b3..c72b0b9 100644 --- a/iree/compiler/Dialect/HAL/Utils/InferCustomKernelsTargetInfoFromParent.cpp +++ b/iree/compiler/Dialect/HAL/Utils/InferCustomKernelsTargetInfoFromParent.cpp
@@ -14,7 +14,7 @@ namespace iree_compiler { LogicalResult InferCustomKernelsTargetInfoFromParent( - FuncOp entryPointFn, CustomKernelsTargetInfo &targetInfo) { + func::FuncOp entryPointFn, CustomKernelsTargetInfo &targetInfo) { // Set the out-value to defaults early so that early returns produce // consistent results and so that we can write simpler code below // (for loop OR-ing booleans, assuming initial 'false' value).
diff --git a/iree/compiler/Dialect/HAL/Utils/InferCustomKernelsTargetInfoFromParent.h b/iree/compiler/Dialect/HAL/Utils/InferCustomKernelsTargetInfoFromParent.h index e5f4aff..ecb415c 100644 --- a/iree/compiler/Dialect/HAL/Utils/InferCustomKernelsTargetInfoFromParent.h +++ b/iree/compiler/Dialect/HAL/Utils/InferCustomKernelsTargetInfoFromParent.h
@@ -12,6 +12,7 @@ #include <cassert> #include "iree/compiler/Utils/CustomKernelsTargetInfo.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/Support/LogicalResult.h" @@ -19,7 +20,7 @@ namespace iree_compiler { LogicalResult InferCustomKernelsTargetInfoFromParent( - FuncOp entryPointFn, CustomKernelsTargetInfo &targetInfo); + func::FuncOp entryPointFn, CustomKernelsTargetInfo &targetInfo); } // namespace iree_compiler } // namespace mlir
diff --git a/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.cpp b/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.cpp index c3fd120..632483f 100644 --- a/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.cpp +++ b/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.cpp
@@ -65,7 +65,7 @@ /// ) LogicalResult updateHALToVMVXEntryFuncOp(FuncOp funcOp, TypeConverter &typeConverter) { - auto originalType = funcOp.getType(); + auto originalType = funcOp.getFunctionType(); if (originalType.getNumInputs() != 0 || originalType.getNumResults() != 0) { return funcOp.emitError() << "exported functions must have no I/O"; } @@ -118,7 +118,7 @@ } // Get the argument to the function corresponding to the workgroup dim. - auto workgroupDim = op->getParentOfType<mlir::FuncOp>().getArgument( + auto workgroupDim = op->getParentOfType<mlir::func::FuncOp>().getArgument( kEntryArgWorkgroupX + dim); rewriter.replaceOp(op, workgroupDim); return success(); @@ -141,7 +141,7 @@ } // Get the argument to the function corresponding to the workgroup dim. - auto workgroupDim = op->getParentOfType<mlir::FuncOp>().getArgument( + auto workgroupDim = op->getParentOfType<mlir::func::FuncOp>().getArgument( kEntryArgWorkgroupSizeX + dim); rewriter.replaceOp(op, workgroupDim); return success(); @@ -164,7 +164,7 @@ } // Get the argument to the function corresponding to the workgroup dim. - auto workgroupDim = op->getParentOfType<mlir::FuncOp>().getArgument( + auto workgroupDim = op->getParentOfType<mlir::func::FuncOp>().getArgument( kEntryArgWorkgroupCountX + dim); rewriter.replaceOp(op, workgroupDim); return success(); @@ -181,8 +181,8 @@ IREE::HAL::InterfaceConstantLoadOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { // Find the vmvx.interface argument to the function. - auto constantsArg = - op->getParentOfType<mlir::FuncOp>().getArgument(kEntryArgConstants); + auto constantsArg = op->getParentOfType<mlir::func::FuncOp>().getArgument( + kEntryArgConstants); assert(constantsArg && "entry point not conforming to requirements"); auto constantType = constantsArg.getType().cast<MemRefType>().getElementType(); @@ -209,8 +209,8 @@ IREE::HAL::InterfaceBindingSubspanOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { // Find the vmvx.interface argument to the function. - auto bindingsArg = - op->getParentOfType<mlir::FuncOp>().getArgument(kEntryArgBindings); + auto bindingsArg = op->getParentOfType<mlir::func::FuncOp>().getArgument( + kEntryArgBindings); assert(bindingsArg && bindingsArg.getType().isa<IREE::Util::ListType>() && "entry point not conforming to requirements");
diff --git a/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.h b/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.h index 8cfc8ed..246fb0f 100644 --- a/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.h +++ b/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.h
@@ -7,6 +7,7 @@ #ifndef IREE_COMPILER_DIALECT_MODULES_VMVX_CONVERSION_HALTOVMVX_CONVERTHALTOVMVX_H_ #define IREE_COMPILER_DIALECT_MODULES_VMVX_CONVERSION_HALTOVMVX_CONVERTHALTOVMVX_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" @@ -15,7 +16,7 @@ // Converts a `() -> ()` function to the calling convention used by VMVX for // passing in bindings, constants, and workgroup parameters. -LogicalResult updateHALToVMVXEntryFuncOp(FuncOp funcOp, +LogicalResult updateHALToVMVXEntryFuncOp(func::FuncOp funcOp, TypeConverter &typeConverter); // Populates conversion patterns from the IREE HAL dialect interface to the
diff --git a/iree/compiler/Dialect/Stream/Analysis/ResourceUsage.cpp b/iree/compiler/Dialect/Stream/Analysis/ResourceUsage.cpp index f67d592..ba6600c 100644 --- a/iree/compiler/Dialect/Stream/Analysis/ResourceUsage.cpp +++ b/iree/compiler/Dialect/Stream/Analysis/ResourceUsage.cpp
@@ -605,7 +605,7 @@ ResourceUsageAnalysis::ResourceUsageAnalysis(Operation *rootOp) : explorer(rootOp, TraversalAction::SHALLOW), solver(explorer, allocator) { explorer.setOpAction<IREE::Util::InitializerOp>(TraversalAction::RECURSE); - explorer.setOpAction<mlir::FuncOp>(TraversalAction::RECURSE); + explorer.setOpAction<mlir::func::FuncOp>(TraversalAction::RECURSE); explorer.setDialectAction<IREE::Stream::StreamDialect>( TraversalAction::RECURSE); // Ignore the contents of executables (linalg goo, etc).
diff --git a/iree/compiler/Dialect/Stream/Conversion/FlowToStream/ConvertFlowToStream.cpp b/iree/compiler/Dialect/Stream/Conversion/FlowToStream/ConvertFlowToStream.cpp index e1cb883..830327f 100644 --- a/iree/compiler/Dialect/Stream/Conversion/FlowToStream/ConvertFlowToStream.cpp +++ b/iree/compiler/Dialect/Stream/Conversion/FlowToStream/ConvertFlowToStream.cpp
@@ -391,7 +391,7 @@ // Update the entry point signatures in the module. // Dispatch tensor arguments become bindings and all others are preserved as // adaptor. Note that we only touch public (exported) functions. - for (auto funcOp : moduleOp.getOps<mlir::FuncOp>()) { + for (auto funcOp : moduleOp.getOps<mlir::func::FuncOp>()) { if (!funcOp.isPublic()) continue; SmallVector<Type> newTypes;
diff --git a/iree/compiler/Dialect/Stream/Conversion/StandardToStream/ConvertStructuralOps.cpp b/iree/compiler/Dialect/Stream/Conversion/StandardToStream/ConvertStructuralOps.cpp index 30e9024..0359101 100644 --- a/iree/compiler/Dialect/Stream/Conversion/StandardToStream/ConvertStructuralOps.cpp +++ b/iree/compiler/Dialect/Stream/Conversion/StandardToStream/ConvertStructuralOps.cpp
@@ -22,16 +22,17 @@ namespace iree_compiler { namespace { -struct FuncOpSignatureConversion : public OpConversionPattern<mlir::FuncOp> { +struct FuncOpSignatureConversion + : public OpConversionPattern<mlir::func::FuncOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - mlir::FuncOp funcOp, OpAdaptor adaptor, + mlir::func::FuncOp funcOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { auto &typeConverter = *getTypeConverter(); // Convert the input signature types. // TODO(benvanik): dynamic shapes by passing in tensor dynamic dims. - auto originalType = funcOp.getType(); + auto originalType = funcOp.getFunctionType(); TypeConverter::SignatureConversion newSignature( originalType.getNumInputs()); for (auto argType : llvm::enumerate(originalType.getInputs())) { @@ -221,10 +222,11 @@ // dynamic legality checker to force any ops using such types to run through // our patterns. - conversionTarget.addDynamicallyLegalOp<mlir::FuncOp>([&](mlir::FuncOp op) { - return typeConverter.isSignatureLegal(op.getType()) && - typeConverter.isLegal(&op.getBody()); - }); + conversionTarget.addDynamicallyLegalOp<mlir::func::FuncOp>( + [&](mlir::func::FuncOp op) { + return typeConverter.isSignatureLegal(op.getFunctionType()) && + typeConverter.isLegal(&op.getBody()); + }); conversionTarget.addDynamicallyLegalOp<mlir::func::CallOp>( [&](mlir::func::CallOp op) { return llvm::all_of(
diff --git a/iree/compiler/Dialect/Stream/IR/StreamOps.cpp b/iree/compiler/Dialect/Stream/IR/StreamOps.cpp index 35576d5..ec5e8a7 100644 --- a/iree/compiler/Dialect/Stream/IR/StreamOps.cpp +++ b/iree/compiler/Dialect/Stream/IR/StreamOps.cpp
@@ -1748,7 +1748,8 @@ // This is sloppy because the function has interleaved bindings and operands; // if we had our own op we could just reuse the map we have for operands. // static -SmallVector<unsigned> CmdDispatchOp::makeOperandToArgMap(mlir::FuncOp funcOp) { +SmallVector<unsigned> CmdDispatchOp::makeOperandToArgMap( + mlir::func::FuncOp funcOp) { unsigned operandCount = llvm::count_if( funcOp.getArgumentTypes(), [](Type type) { return !type.isa<IREE::Stream::BindingType>(); }); @@ -1765,7 +1766,8 @@ } // static -SmallVector<unsigned> CmdDispatchOp::makeResourceToArgMap(mlir::FuncOp funcOp) { +SmallVector<unsigned> CmdDispatchOp::makeResourceToArgMap( + mlir::func::FuncOp funcOp) { unsigned operandCount = llvm::count_if( funcOp.getArgumentTypes(), [](Type type) { return type.isa<IREE::Stream::BindingType>(); }); @@ -2023,11 +2025,11 @@ builder.getStringAttr(sym_name), function_ref); } -::mlir::FuncOp ExecutableExportOp::getFunctionRef() { +::mlir::func::FuncOp ExecutableExportOp::getFunctionRef() { auto executableOp = this->getOperation()->getParentOfType<IREE::Stream::ExecutableOp>(); if (!executableOp) return {}; - return executableOp.getInnerModule().lookupSymbol<::mlir::FuncOp>( + return executableOp.getInnerModule().lookupSymbol<::mlir::func::FuncOp>( function_ref()); }
diff --git a/iree/compiler/Dialect/Stream/IR/StreamOps.h b/iree/compiler/Dialect/Stream/IR/StreamOps.h index f781525..8b17a0f 100644 --- a/iree/compiler/Dialect/Stream/IR/StreamOps.h +++ b/iree/compiler/Dialect/Stream/IR/StreamOps.h
@@ -15,6 +15,7 @@ #include "iree/compiler/Dialect/Util/IR/UtilOps.h" #include "iree/compiler/Dialect/Util/IR/UtilTraits.h" #include "iree/compiler/Dialect/Util/IR/UtilTypes.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/BuiltinTypes.h"
diff --git a/iree/compiler/Dialect/Stream/IR/StreamOps.td b/iree/compiler/Dialect/Stream/IR/StreamOps.td index a4d5c73..605f575 100644 --- a/iree/compiler/Dialect/Stream/IR/StreamOps.td +++ b/iree/compiler/Dialect/Stream/IR/StreamOps.td
@@ -2328,9 +2328,9 @@ Value getResultSize(unsigned idx) { return {}; } // Builds a map of operand index to argument index. - static SmallVector<unsigned> makeOperandToArgMap(mlir::FuncOp funcOp); + static SmallVector<unsigned> makeOperandToArgMap(mlir::func::FuncOp funcOp); // Builds a map of resource to argument index of the corresponding binding. - static SmallVector<unsigned> makeResourceToArgMap(mlir::FuncOp funcOp); + static SmallVector<unsigned> makeResourceToArgMap(mlir::func::FuncOp funcOp); }]; let hasVerifier = 1; @@ -2805,7 +2805,7 @@ ]; let extraClassDeclaration = [{ - ::mlir::FuncOp getFunctionRef(); + ::mlir::func::FuncOp getFunctionRef(); }]; }
diff --git a/iree/compiler/Dialect/Stream/Transforms/AnnotateDispatchArguments.cpp b/iree/compiler/Dialect/Stream/Transforms/AnnotateDispatchArguments.cpp index 08b65e9..05c3116 100644 --- a/iree/compiler/Dialect/Stream/Transforms/AnnotateDispatchArguments.cpp +++ b/iree/compiler/Dialect/Stream/Transforms/AnnotateDispatchArguments.cpp
@@ -324,7 +324,7 @@ : explorer(rootOp, TraversalAction::SHALLOW), solver(explorer, allocator) { explorer.setOpAction<IREE::Util::InitializerOp>(TraversalAction::RECURSE); - explorer.setOpAction<mlir::FuncOp>(TraversalAction::RECURSE); + explorer.setOpAction<mlir::func::FuncOp>(TraversalAction::RECURSE); explorer.setDialectAction<IREE::Stream::StreamDialect>( TraversalAction::RECURSE); // Ignore the contents of executables (linalg goo, etc).
diff --git a/iree/compiler/Dialect/Stream/Transforms/DumpStatistics.cpp b/iree/compiler/Dialect/Stream/Transforms/DumpStatistics.cpp index 10e4acc..0584291 100644 --- a/iree/compiler/Dialect/Stream/Transforms/DumpStatistics.cpp +++ b/iree/compiler/Dialect/Stream/Transforms/DumpStatistics.cpp
@@ -43,7 +43,7 @@ // stream.executable ops mapped by name. llvm::MapVector<StringRef, IREE::Stream::ExecutableOp> executableOps; // stream.executable exported function -> dispatches to it. - llvm::MapVector<mlir::FuncOp, SmallVector<IREE::Stream::CmdDispatchOp>> + llvm::MapVector<mlir::func::FuncOp, SmallVector<IREE::Stream::CmdDispatchOp>> exportDispatchOps; // TODO(benvanik): resource allocations.
diff --git a/iree/compiler/Dialect/Stream/Transforms/ElideAsyncCopies.cpp b/iree/compiler/Dialect/Stream/Transforms/ElideAsyncCopies.cpp index 27e09e1..5331b93 100644 --- a/iree/compiler/Dialect/Stream/Transforms/ElideAsyncCopies.cpp +++ b/iree/compiler/Dialect/Stream/Transforms/ElideAsyncCopies.cpp
@@ -290,7 +290,7 @@ : explorer(rootOp, TraversalAction::SHALLOW), solver(explorer, allocator) { explorer.setOpAction<IREE::Util::InitializerOp>(TraversalAction::RECURSE); - explorer.setOpAction<mlir::FuncOp>(TraversalAction::RECURSE); + explorer.setOpAction<mlir::func::FuncOp>(TraversalAction::RECURSE); explorer.setDialectAction<IREE::Stream::StreamDialect>( TraversalAction::RECURSE); // Ignore the contents of executables (linalg goo, etc).
diff --git a/iree/compiler/Dialect/Stream/Transforms/FoldUniformOperands.cpp b/iree/compiler/Dialect/Stream/Transforms/FoldUniformOperands.cpp index b4ea95f..37e1b26 100644 --- a/iree/compiler/Dialect/Stream/Transforms/FoldUniformOperands.cpp +++ b/iree/compiler/Dialect/Stream/Transforms/FoldUniformOperands.cpp
@@ -43,7 +43,7 @@ // stream.cmd.dispatch @foo(%0, %1 : index, index) // + deduped arguments in the executable static void deduplicateOperands( - mlir::FuncOp funcOp, + mlir::func::FuncOp funcOp, SmallVector<IREE::Stream::CmdDispatchOp> &dispatchOps) { auto &entryBlock = funcOp.front(); auto anyDispatchOp = dispatchOps.front(); @@ -115,7 +115,7 @@ } LLVM_DEBUG({ - llvm::dbgs() << "deduplicateOperands for " << funcOp.sym_name() << "\n"; + llvm::dbgs() << "deduplicateOperands for " << funcOp.getSymName() << "\n"; llvm::dbgs() << " dead operands: "; llvm::interleaveComma(deadOperandsMap.set_bits(), llvm::dbgs()); llvm::dbgs() << "\n"; @@ -165,7 +165,7 @@ // stream.cmd.dispatch @foo(%c101 : index) // + inlined %c1 in the executable static void inlineUniformConstants( - mlir::FuncOp funcOp, + mlir::func::FuncOp funcOp, SmallVector<IREE::Stream::CmdDispatchOp> &dispatchOps) { auto &entryBlock = funcOp.front(); auto anyDispatchOp = dispatchOps.front(); @@ -204,7 +204,8 @@ } LLVM_DEBUG({ - llvm::dbgs() << "inlineUniformConstants for " << funcOp.sym_name() << "\n"; + llvm::dbgs() << "inlineUniformConstants for " << funcOp.getSymName() + << "\n"; for (unsigned i = 0; i < operandValues.size(); ++i) { if (!operandValues[i].hasValue()) continue; llvm::dbgs() << " operand " << i << " = " << operandValues[i].getValue()
diff --git a/iree/compiler/Dialect/Stream/Transforms/FuseDispatchBindings.cpp b/iree/compiler/Dialect/Stream/Transforms/FuseDispatchBindings.cpp index e55e705..e7fd525 100644 --- a/iree/compiler/Dialect/Stream/Transforms/FuseDispatchBindings.cpp +++ b/iree/compiler/Dialect/Stream/Transforms/FuseDispatchBindings.cpp
@@ -172,7 +172,7 @@ // Updates an executable function to use the new bindings. static void updateExecutableSignature(IREE::Stream::ExecutableOp executableOp, IREE::Stream::ExecutableExportOp exportOp, - mlir::FuncOp funcOp, + mlir::func::FuncOp funcOp, ArrayRef<Binding> bindings) { auto &entryBlock = funcOp.front();
diff --git a/iree/compiler/Dialect/Stream/Transforms/Passes.cpp b/iree/compiler/Dialect/Stream/Transforms/Passes.cpp index 7ecf05a..8cbbcad 100644 --- a/iree/compiler/Dialect/Stream/Transforms/Passes.cpp +++ b/iree/compiler/Dialect/Stream/Transforms/Passes.cpp
@@ -31,7 +31,7 @@ // redundant store-loads are removed. passManager.addNestedPass<IREE::Util::InitializerOp>( IREE::Util::createSimplifyGlobalAccessesPass()); - passManager.addNestedPass<mlir::FuncOp>( + passManager.addNestedPass<mlir::func::FuncOp>( IREE::Util::createSimplifyGlobalAccessesPass()); // Cleanup and canonicalization of util.global (and other util ops). @@ -109,7 +109,7 @@ // affinity/configuration assigned during placement. passManager.addNestedPass<IREE::Util::InitializerOp>( IREE::Stream::createEncodeHostTensorsPass()); - passManager.addNestedPass<mlir::FuncOp>( + passManager.addNestedPass<mlir::func::FuncOp>( IREE::Stream::createEncodeHostTensorsPass()); passManager.addNestedPass<IREE::Stream::ExecutableOp>( IREE::Stream::createEncodeDeviceTensorsPass()); @@ -127,7 +127,7 @@ // (though it's critical enough that it is not optional). passManager.addNestedPass<IREE::Util::InitializerOp>( IREE::Stream::createMaterializeCopyOnWritePass()); - passManager.addNestedPass<mlir::FuncOp>( + passManager.addNestedPass<mlir::func::FuncOp>( IREE::Stream::createMaterializeCopyOnWritePass()); passManager.addPass(IREE::Stream::createElideAsyncCopiesPass()); @@ -144,13 +144,13 @@ // Combine async work into execution regions. passManager.addNestedPass<IREE::Util::InitializerOp>( IREE::Stream::createScheduleExecutionPass()); - passManager.addNestedPass<mlir::FuncOp>( + passManager.addNestedPass<mlir::func::FuncOp>( IREE::Stream::createScheduleExecutionPass()); // Group concurrently executable work into waves. passManager.addNestedPass<IREE::Util::InitializerOp>( IREE::Stream::createScheduleConcurrencyPass()); - passManager.addNestedPass<mlir::FuncOp>( + passManager.addNestedPass<mlir::func::FuncOp>( IREE::Stream::createScheduleConcurrencyPass()); // Materialize timepoints across the entire module. This simplifies scheduling @@ -175,7 +175,7 @@ // lifetime allocations. passManager.addNestedPass<IREE::Util::InitializerOp>( IREE::Stream::createScheduleAllocationPass()); - passManager.addNestedPass<mlir::FuncOp>( + passManager.addNestedPass<mlir::func::FuncOp>( IREE::Stream::createScheduleAllocationPass()); // TODO(benvanik): passes to convert alloc to alloca and thread through @@ -187,13 +187,13 @@ // buffers and upload logic. passManager.addNestedPass<IREE::Util::InitializerOp>( IREE::Stream::createPackConstantsPass()); - passManager.addNestedPass<mlir::FuncOp>( + passManager.addNestedPass<mlir::func::FuncOp>( IREE::Stream::createPackConstantsPass()); // Pack fused allocations based on lifetime. passManager.addNestedPass<IREE::Util::InitializerOp>( IREE::Stream::createPackAllocationsPass()); - passManager.addNestedPass<mlir::FuncOp>( + passManager.addNestedPass<mlir::func::FuncOp>( IREE::Stream::createPackAllocationsPass()); // Layout packed slices to emit the arithmetic required for all resource @@ -201,7 +201,7 @@ // below. passManager.addNestedPass<IREE::Util::InitializerOp>( IREE::Stream::createLayoutSlicesPass()); - passManager.addNestedPass<mlir::FuncOp>( + passManager.addNestedPass<mlir::func::FuncOp>( IREE::Stream::createLayoutSlicesPass()); // Propagate subviews throughout the program to unify resource storage access.
diff --git a/iree/compiler/Dialect/Stream/Transforms/PropagateSubviews.cpp b/iree/compiler/Dialect/Stream/Transforms/PropagateSubviews.cpp index 24c15fd..8caba29 100644 --- a/iree/compiler/Dialect/Stream/Transforms/PropagateSubviews.cpp +++ b/iree/compiler/Dialect/Stream/Transforms/PropagateSubviews.cpp
@@ -357,9 +357,9 @@ // -> // func @foo(%0: !stream.resource, %sz: index, %o: index, %l: index) { // %1 = stream.resource.subview %0[%o] : {%sz} -> {%l} -static void expandFuncOp(mlir::FuncOp op, ExpandedGlobalMap &globalMap, +static void expandFuncOp(mlir::func::FuncOp op, ExpandedGlobalMap &globalMap, IndexSet &indexSet, SubviewMap &subviewMap) { - auto oldType = op.getType(); + auto oldType = op.getFunctionType(); auto inputTypes = expandTypes(oldType.getInputs()); auto resultTypes = expandTypes(oldType.getResults()); auto newType = FunctionType::get(op.getContext(), inputTypes, resultTypes); @@ -484,7 +484,7 @@ expandGlobalStoreOp(storeOp, globalMap, indexSet, subviewMap); } else if (auto initializerOp = dyn_cast<IREE::Util::InitializerOp>(op)) { expandInitializerOp(initializerOp, globalMap, indexSet, subviewMap); - } else if (auto funcOp = dyn_cast<mlir::FuncOp>(op)) { + } else if (auto funcOp = dyn_cast<mlir::func::FuncOp>(op)) { expandFuncOp(funcOp, globalMap, indexSet, subviewMap); } else if (auto callOp = dyn_cast<mlir::func::CallOp>(op)) { expandCallOp(callOp, indexSet, subviewMap);
diff --git a/iree/compiler/Dialect/Stream/Transforms/PropagateTimepoints.cpp b/iree/compiler/Dialect/Stream/Transforms/PropagateTimepoints.cpp index d2b0e55..78789fd 100644 --- a/iree/compiler/Dialect/Stream/Transforms/PropagateTimepoints.cpp +++ b/iree/compiler/Dialect/Stream/Transforms/PropagateTimepoints.cpp
@@ -362,9 +362,9 @@ // -> // func @foo(%t: !stream.timepoint, %0: !stream.resource) { // %1 = stream.timepoint.await %t, %0 -static void expandFuncOp(mlir::FuncOp op, ExpandedGlobalMap &globalMap, +static void expandFuncOp(mlir::func::FuncOp op, ExpandedGlobalMap &globalMap, BlockAndValueMapping &resourceTimepointMap) { - auto oldType = op.getType(); + auto oldType = op.getFunctionType(); auto inputTypes = expandTypes(oldType.getInputs()); auto resultTypes = expandTypes(oldType.getResults()); auto newType = FunctionType::get(op.getContext(), inputTypes, resultTypes); @@ -545,7 +545,7 @@ expandGlobalStoreOp(storeOp, globalMap, resourceTimepointMap); } else if (auto initializerOp = dyn_cast<IREE::Util::InitializerOp>(op)) { expandInitializerOp(initializerOp, globalMap, resourceTimepointMap); - } else if (auto funcOp = dyn_cast<mlir::FuncOp>(op)) { + } else if (auto funcOp = dyn_cast<mlir::func::FuncOp>(op)) { expandFuncOp(funcOp, globalMap, resourceTimepointMap); } else if (auto callOp = dyn_cast<mlir::func::CallOp>(op)) { expandCallOp(callOp, resourceTimepointMap);
diff --git a/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp b/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp index d05c4cf..3d26f94 100644 --- a/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp +++ b/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp
@@ -205,15 +205,15 @@ // Applies usage analysis results to an MLIR function. // All resource arguments and results, block arguments, and nested operations // will have their lifetime specified. -struct ApplyFuncOp : public UsageRefinementPattern<mlir::FuncOp> { - using UsageRefinementPattern<mlir::FuncOp>::UsageRefinementPattern; - LogicalResult matchAndRewrite(mlir::FuncOp op, +struct ApplyFuncOp : public UsageRefinementPattern<mlir::func::FuncOp> { + using UsageRefinementPattern<mlir::func::FuncOp>::UsageRefinementPattern; + LogicalResult matchAndRewrite(mlir::func::FuncOp op, PatternRewriter &rewriter) const override { bool didChange = false; // Arguments: SmallVector<Type> newInputs; - for (auto inputType : llvm::enumerate(op.getType().getInputs())) { + for (auto inputType : llvm::enumerate(op.getFunctionType().getInputs())) { auto oldType = inputType.value().dyn_cast<IREE::Stream::ResourceType>(); if (!oldType) { newInputs.push_back(inputType.value()); @@ -232,7 +232,7 @@ // Results: SmallVector<Type> newOutputs; auto anyReturnOp = *op.getOps<mlir::func::ReturnOp>().begin(); - for (auto outputType : llvm::enumerate(op.getType().getResults())) { + for (auto outputType : llvm::enumerate(op.getFunctionType().getResults())) { auto oldType = outputType.value().dyn_cast<IREE::Stream::ResourceType>(); if (!oldType) { newOutputs.push_back(outputType.value()); @@ -248,7 +248,7 @@ } } auto newFuncType = rewriter.getFunctionType(newInputs, newOutputs); - if (op.getType() != newFuncType) { + if (op.getFunctionType() != newFuncType) { op.setType(newFuncType); didChange = true; }
diff --git a/iree/compiler/Dialect/Stream/Transforms/SpecializeDispatches.cpp b/iree/compiler/Dialect/Stream/Transforms/SpecializeDispatches.cpp index 75bc1bf..abfcad4 100644 --- a/iree/compiler/Dialect/Stream/Transforms/SpecializeDispatches.cpp +++ b/iree/compiler/Dialect/Stream/Transforms/SpecializeDispatches.cpp
@@ -55,7 +55,7 @@ // Each dispatch gets a row in the table that can be selected based on the // dispatch ordinal. static ConstantTable buildConstantTable( - mlir::FuncOp funcOp, + mlir::func::FuncOp funcOp, SmallVector<IREE::Stream::CmdDispatchOp> &dispatchOps) { auto anyDispatchOp = dispatchOps.front(); unsigned operandCount = anyDispatchOp.operands().size(); @@ -161,7 +161,7 @@ // // TODO(benvanik): maybe a dedicated lookup table op to make further combining // easier to do in a backend-generic way. -static void insertConstantTableLookup(mlir::FuncOp funcOp, +static void insertConstantTableLookup(mlir::func::FuncOp funcOp, ConstantTable &constantTable) { auto &entryBlock = funcOp.front(); auto operandToArgMap =
diff --git a/iree/compiler/Dialect/Stream/Transforms/test/fold_globals.mlir b/iree/compiler/Dialect/Stream/Transforms/test/fold_globals.mlir index b58d5ba..fefaa46 100644 --- a/iree/compiler/Dialect/Stream/Transforms/test/fold_globals.mlir +++ b/iree/compiler/Dialect/Stream/Transforms/test/fold_globals.mlir
@@ -6,13 +6,13 @@ // CHECK: util.global public mutable @uniformConstants = #stream.timepoint<immediate> util.global public mutable @uniformConstants : !stream.timepoint -builtin.func @foo() { +func.func @foo() { %timepoint = stream.timepoint.immediate => !stream.timepoint // CHECK-NOT: util.global.store util.global.store %timepoint, @uniformConstants : !stream.timepoint return } -builtin.func @bar() { +func.func @bar() { %timepoint = stream.timepoint.immediate => !stream.timepoint // CHECK-NOT: util.global.store util.global.store %timepoint, @uniformConstants : !stream.timepoint @@ -23,7 +23,7 @@ // CHECK-NOT: @immutable util.global private @immutable = #stream.timepoint<immediate> : !stream.timepoint -builtin.func @foo() -> !stream.timepoint { +func.func @foo() -> !stream.timepoint { // CHECK-NOT: util.global.load @immutable // CHECK: %[[IMMEDIATE:.+]] = stream.timepoint.immediate => !stream.timepoint %0 = util.global.load @immutable : !stream.timepoint
diff --git a/iree/compiler/Dialect/Stream/Transforms/test/layout_slices.mlir b/iree/compiler/Dialect/Stream/Transforms/test/layout_slices.mlir index 26db0ee..0b8bb31 100644 --- a/iree/compiler/Dialect/Stream/Transforms/test/layout_slices.mlir +++ b/iree/compiler/Dialect/Stream/Transforms/test/layout_slices.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='builtin.func(iree-stream-layout-slices)' -cse %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='func.func(iree-stream-layout-slices)' -cse %s | FileCheck %s #layoutStaticConfig = #stream.resource_config<{ max_allocation_size = 1073741824,
diff --git a/iree/compiler/Dialect/Stream/Transforms/test/materialize_copy_on_write.mlir b/iree/compiler/Dialect/Stream/Transforms/test/materialize_copy_on_write.mlir index 12e02d8..3242934 100644 --- a/iree/compiler/Dialect/Stream/Transforms/test/materialize_copy_on_write.mlir +++ b/iree/compiler/Dialect/Stream/Transforms/test/materialize_copy_on_write.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='builtin.func(iree-stream-materialize-copy-on-write)' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='func.func(iree-stream-materialize-copy-on-write)' %s | FileCheck %s // Tests that block arguments (including function arguments) are always cloned. // Until a whole-program analysis runs we don't know their semantics.
diff --git a/iree/compiler/Dialect/Stream/Transforms/test/pack_allocations.mlir b/iree/compiler/Dialect/Stream/Transforms/test/pack_allocations.mlir index c14d3b9..8e2550d 100644 --- a/iree/compiler/Dialect/Stream/Transforms/test/pack_allocations.mlir +++ b/iree/compiler/Dialect/Stream/Transforms/test/pack_allocations.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='builtin.func(iree-stream-pack-allocations)' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='func.func(iree-stream-pack-allocations)' %s | FileCheck %s // CHECK-LABEL: @packAllocations // CHECK-SAME: (%[[SIZE_A:.+]]: index, %[[SIZE_B:.+]]: index)
diff --git a/iree/compiler/Dialect/Stream/Transforms/test/pack_constants.mlir b/iree/compiler/Dialect/Stream/Transforms/test/pack_constants.mlir index 6c91d54..e0a77f9 100644 --- a/iree/compiler/Dialect/Stream/Transforms/test/pack_constants.mlir +++ b/iree/compiler/Dialect/Stream/Transforms/test/pack_constants.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='builtin.func(iree-stream-pack-constants)' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='func.func(iree-stream-pack-constants)' %s | FileCheck %s // This is a high level test of the structure emitted by the pass. // Subsequent tests focus on individual components.
diff --git a/iree/compiler/Dialect/Stream/Transforms/test/schedule_allocation.mlir b/iree/compiler/Dialect/Stream/Transforms/test/schedule_allocation.mlir index 00c39f9..32592ee 100644 --- a/iree/compiler/Dialect/Stream/Transforms/test/schedule_allocation.mlir +++ b/iree/compiler/Dialect/Stream/Transforms/test/schedule_allocation.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='builtin.func(iree-stream-schedule-allocation)' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='func.func(iree-stream-schedule-allocation)' %s | FileCheck %s // Tests that async constant ops get extracted into a dedicated constant op // outside of the execution region. This allows us to handle them in various
diff --git a/iree/compiler/Dialect/Stream/Transforms/test/schedule_concurrency.mlir b/iree/compiler/Dialect/Stream/Transforms/test/schedule_concurrency.mlir index 794d510..60a9383 100644 --- a/iree/compiler/Dialect/Stream/Transforms/test/schedule_concurrency.mlir +++ b/iree/compiler/Dialect/Stream/Transforms/test/schedule_concurrency.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline="builtin.func(iree-stream-schedule-concurrency)" %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline="func.func(iree-stream-schedule-concurrency)" %s | FileCheck %s // Tests that when favor=min-peak-memory we assume ops are in an order that // reduces live memory ranges and only optimistically put them in concurrency
diff --git a/iree/compiler/Dialect/Stream/Transforms/test/schedule_execution.mlir b/iree/compiler/Dialect/Stream/Transforms/test/schedule_execution.mlir index 8524578..02a264b 100644 --- a/iree/compiler/Dialect/Stream/Transforms/test/schedule_execution.mlir +++ b/iree/compiler/Dialect/Stream/Transforms/test/schedule_execution.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline="builtin.func(iree-stream-schedule-execution)" %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline="func.func(iree-stream-schedule-execution)" %s | FileCheck %s // Tests basic partitioning of multiple ops.
diff --git a/iree/compiler/Dialect/Util/IR/UtilDialect.cpp b/iree/compiler/Dialect/Util/IR/UtilDialect.cpp index ea39e8b..db3fd35 100644 --- a/iree/compiler/Dialect/Util/IR/UtilDialect.cpp +++ b/iree/compiler/Dialect/Util/IR/UtilDialect.cpp
@@ -16,6 +16,7 @@ #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/DialectImplementation.h" +#include "mlir/IR/MLIRContext.h" #include "mlir/IR/Matchers.h" #include "mlir/IR/OpDefinition.h" #include "mlir/IR/OpImplementation.h" @@ -159,14 +160,14 @@ OpTy> {}; template <typename OpTy> - static void add(DialectRegistry ®istry) { - registry.addOpInterface<OpTy, ExternalModel<OpTy>>(); + static void add(MLIRContext *ctx) { + OpTy::template attachInterface<ExternalModel<OpTy>>(*ctx); } template <typename OpTy1, typename OpTy2, typename... More> - static void add(DialectRegistry ®istry) { - add<OpTy1>(registry); - add<OpTy2, More...>(registry); + static void add(MLIRContext *ctx) { + add<OpTy1>(ctx); + add<OpTy2, More...>(ctx); } }; @@ -176,10 +177,13 @@ // Must ensure that any dependent dialects are registered. registry.insert<arith::ArithmeticDialect>(); - GenericNumericCastExternalModel::add< - arith::BitcastOp, arith::ExtFOp, arith::ExtUIOp, arith::ExtSIOp, - arith::FPToSIOp, arith::FPToUIOp, arith::IndexCastOp, arith::TruncFOp, - arith::TruncIOp, arith::SIToFPOp, arith::UIToFPOp>(registry); + registry.addExtension(+[](MLIRContext *ctx, + arith::ArithmeticDialect *dialect) { + GenericNumericCastExternalModel::add< + arith::BitcastOp, arith::ExtFOp, arith::ExtUIOp, arith::ExtSIOp, + arith::FPToSIOp, arith::FPToUIOp, arith::IndexCastOp, arith::TruncFOp, + arith::TruncIOp, arith::SIToFPOp, arith::UIToFPOp>(ctx); + }); } } // namespace Util
diff --git a/iree/compiler/Dialect/Util/IR/UtilOps.cpp b/iree/compiler/Dialect/Util/IR/UtilOps.cpp index 45211fd..31e5233 100644 --- a/iree/compiler/Dialect/Util/IR/UtilOps.cpp +++ b/iree/compiler/Dialect/Util/IR/UtilOps.cpp
@@ -688,15 +688,15 @@ void InitializerOp::build(OpBuilder &builder, OperationState &result, ArrayRef<NamedAttribute> attrs) { - result.addAttribute( - "type", TypeAttr::get(FunctionType::get(builder.getContext(), {}, {}))); + result.addAttribute("function_type", TypeAttr::get(FunctionType::get( + builder.getContext(), {}, {}))); result.addRegion(); result.attributes.append(attrs.begin(), attrs.end()); } ParseResult InitializerOp::parse(OpAsmParser &parser, OperationState &result) { - result.addAttribute( - "type", TypeAttr::get(FunctionType::get(result.getContext(), {}, {}))); + result.addAttribute("function_type", TypeAttr::get(FunctionType::get( + result.getContext(), {}, {}))); if (parser.parseOptionalAttrDictWithKeyword(result.attributes)) { return failure(); } @@ -709,7 +709,8 @@ void InitializerOp::print(OpAsmPrinter &p) { Operation *op = getOperation(); - p.printOptionalAttrDictWithKeyword(op->getAttrs(), /*elidedAttrs=*/{"type"}); + p.printOptionalAttrDictWithKeyword(op->getAttrs(), + /*elidedAttrs=*/{"function_type"}); p << " "; p.printRegion(body()); }
diff --git a/iree/compiler/Dialect/Util/IR/UtilOps.h b/iree/compiler/Dialect/Util/IR/UtilOps.h index f9c1f0f..0793c49 100644 --- a/iree/compiler/Dialect/Util/IR/UtilOps.h +++ b/iree/compiler/Dialect/Util/IR/UtilOps.h
@@ -12,9 +12,11 @@ #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/Dialect.h" +#include "mlir/IR/FunctionInterfaces.h" #include "mlir/IR/OpDefinition.h" #include "mlir/IR/OpImplementation.h" #include "mlir/IR/SymbolTable.h" +#include "mlir/Interfaces/CallInterfaces.h" #include "mlir/Interfaces/ControlFlowInterfaces.h" #include "mlir/Interfaces/InferTypeOpInterface.h" #include "mlir/Interfaces/SideEffectInterfaces.h"
diff --git a/iree/compiler/Dialect/Util/IR/UtilOps.td b/iree/compiler/Dialect/Util/IR/UtilOps.td index 08f6782..67a2185 100644 --- a/iree/compiler/Dialect/Util/IR/UtilOps.td +++ b/iree/compiler/Dialect/Util/IR/UtilOps.td
@@ -377,7 +377,7 @@ }]; let arguments = (ins - TypeAttr:$type + TypeAttr:$function_type ); let regions = (region AnyRegion:$body); @@ -395,8 +395,10 @@ Block *addEntryBlock(); Block *addBlock(); - FunctionType getType() { - return getTypeAttr().getValue().cast<FunctionType>(); + FunctionType getFunctionType() { + return getOperation()->getAttrOfType<TypeAttr>(getTypeAttrName()) + .getValue() + .cast<FunctionType>(); } /// Returns the argument types of this function.
diff --git a/iree/compiler/Dialect/Util/Transforms/test/combine_initializers.mlir b/iree/compiler/Dialect/Util/Transforms/test/combine_initializers.mlir index b61a022..53fada5 100644 --- a/iree/compiler/Dialect/Util/Transforms/test/combine_initializers.mlir +++ b/iree/compiler/Dialect/Util/Transforms/test/combine_initializers.mlir
@@ -2,7 +2,7 @@ // Tests that multiple initializers are combined in their module order. -builtin.func private @extern() -> index +func.func private @extern() -> index // CHECK: util.global private mutable @global0 : index util.global private mutable @global0 : index @@ -32,7 +32,7 @@ // CHECK-NEXT: util.initializer.return // CHECK-LABEL: @orderedCombining -builtin.func @orderedCombining(%arg0: index) -> (index, index, index) { +func.func @orderedCombining(%arg0: index) -> (index, index, index) { util.global.store %arg0, @global0 : index %value0 = util.global.load @global0 : index %value1 = util.global.load @global1 : index
diff --git a/iree/compiler/Dialect/Util/Transforms/test/fold_globals.mlir b/iree/compiler/Dialect/Util/Transforms/test/fold_globals.mlir index fd1441a..c86c8d6 100644 --- a/iree/compiler/Dialect/Util/Transforms/test/fold_globals.mlir +++ b/iree/compiler/Dialect/Util/Transforms/test/fold_globals.mlir
@@ -2,13 +2,13 @@ // CHECK: util.global public mutable @uniformConstants = 5 : index util.global public mutable @uniformConstants : index -builtin.func @foo() { +func.func @foo() { %c5 = arith.constant 5 : index // CHECK-NOT: util.global.store %c5, @uniformConstants : index util.global.store %c5, @uniformConstants : index return } -builtin.func @bar() { +func.func @bar() { %c5 = arith.constant 5 : index // CHECK-NOT: util.global.store %c5, @uniformConstants : index util.global.store %c5, @uniformConstants : index @@ -19,13 +19,13 @@ // CHECK: util.global public mutable @nonuniformConstants : index util.global public mutable @nonuniformConstants : index -builtin.func @foo() { +func.func @foo() { %c5 = arith.constant 5 : index // CHECK: util.global.store %c5, @nonuniformConstants : index util.global.store %c5, @nonuniformConstants : index return } -builtin.func @bar() { +func.func @bar() { %c6 = arith.constant 6 : index // CHECK: util.global.store %c6, @nonuniformConstants : index util.global.store %c6, @nonuniformConstants : index @@ -38,7 +38,7 @@ util.global private mutable @chained0 : index // CHECK-NOT: util.global private mutable @chained1 : index util.global private mutable @chained1 : index -builtin.func @foo() -> index { +func.func @foo() -> index { // CHECK: %[[VALUE:.+]] = util.global.load @chained0 : index %0 = util.global.load @chained0 : index // CHECK-NOT: util.global.store @@ -53,14 +53,14 @@ util.global public mutable @unchained0 : index // CHECK: util.global public mutable @unchained1 : index util.global public mutable @unchained1 : index -builtin.func @foo() { +func.func @foo() { // CHECK: %[[VALUE:.+]] = util.global.load @unchained0 : index %0 = util.global.load @unchained0 : index // CHECK: util.global.store %[[VALUE]], @unchained1 : index util.global.store %0, @unchained1 : index return } -builtin.func @bar(%arg0: index) { +func.func @bar(%arg0: index) { // CHECK: util.global.store %arg0, @unchained1 : index util.global.store %arg0, @unchained1 : index return @@ -83,7 +83,7 @@ util.global.store %c6, @immutable1 : index util.initializer.return } -builtin.func @foo(%arg0: index) -> (index, index, index) { +func.func @foo(%arg0: index) -> (index, index, index) { // CHECK-DAG: %[[C5:.+]] = arith.constant 5 %0 = util.global.load @immutable0 : index // CHECK-DAG: %[[C6:.+]] = arith.constant 6 @@ -102,7 +102,7 @@ util.global private mutable @used0 = 5 : index // CHECK: util.global private mutable @used1 : index util.global private mutable @used1 : index -builtin.func @foo(%arg0: index, %arg1: index) -> (index, index) { +func.func @foo(%arg0: index, %arg1: index) -> (index, index) { // CHECK: %[[VALUE0:.+]] = util.global.load @used0 : index %0 = util.global.load @used0 : index // CHECK: %[[VALUE1:.+]] = util.global.load @used1 : index @@ -134,7 +134,7 @@ util.global private @dupeCst0 {noinline} = 5 : index // CHECK-NOT: util.global private @dupeCst1 util.global private @dupeCst1 {noinline} = 5 : index -builtin.func @foo() -> (index, index) { +func.func @foo() -> (index, index) { // CHECK-DAG: %[[VALUE0:.+]] = util.global.load @dupeCst0 %0 = util.global.load @dupeCst0 : index // CHECK-DAG: %[[VALUE1:.+]] = util.global.load @dupeCst0 @@ -155,7 +155,7 @@ util.global.store %c7, @nondupeCst1 : index util.initializer.return } -builtin.func @foo() -> (index, index) { +func.func @foo() -> (index, index) { // CHECK-DAG: %[[C6:.+]] = arith.constant 6 : index %0 = util.global.load @nondupeCst0 : index // CHECK-DAG: %[[C7:.+]] = arith.constant 7 : index
diff --git a/iree/compiler/Dialect/Util/Transforms/test/fuse_globals.mlir b/iree/compiler/Dialect/Util/Transforms/test/fuse_globals.mlir index e3e4767..2fff5b5 100644 --- a/iree/compiler/Dialect/Util/Transforms/test/fuse_globals.mlir +++ b/iree/compiler/Dialect/Util/Transforms/test/fuse_globals.mlir
@@ -3,7 +3,7 @@ // CHECK: util.global private mutable @fusable0 : index util.global private mutable @fusable0 : index util.global private mutable @fusable1 : index -builtin.func @foo(%arg0: index) -> (index, index) { +func.func @foo(%arg0: index) -> (index, index) { // CHECK: util.global.store %arg0, @fusable0 util.global.store %arg0, @fusable0 : index // CHECK-NOT: util.global.store %arg0, @fusable1 @@ -24,7 +24,7 @@ util.global private mutable @unfusable0 : index // CHECK: util.global private mutable @unfusable1 : index util.global private mutable @unfusable1 : index -builtin.func @foo(%arg0: index) -> (index, index) { +func.func @foo(%arg0: index) -> (index, index) { // CHECK: util.global.store %arg0, @unfusable0 : index util.global.store %arg0, @unfusable0 : index // CHECK: util.global.store %arg0, @unfusable1 : index @@ -36,7 +36,7 @@ // CHECK: return %[[VALUE0]], %[[VALUE1]] return %0, %1 : index, index } -builtin.func @bar(%arg0: index) { +func.func @bar(%arg0: index) { util.global.store %arg0, @unfusable0 : index return } @@ -54,7 +54,7 @@ util.global private mutable @unfusableInit0 = 5 : index // CHECK: util.global private mutable @unfusableInit1 = 6 : index util.global private mutable @unfusableInit1 = 6 : index -builtin.func @foo(%arg0: index) -> (index, index) { +func.func @foo(%arg0: index) -> (index, index) { // CHECK: util.global.store %arg0, @unfusableInit0 util.global.store %arg0, @unfusableInit0 : index // CHECK: util.global.store %arg0, @unfusableInit1 @@ -73,12 +73,12 @@ util.global private mutable @unfusableDivergent0 : index // CHECK: util.global private mutable @unfusableDivergent1 util.global private mutable @unfusableDivergent1 : index -builtin.func @fn_a(%arg0: index) { +func.func @fn_a(%arg0: index) { util.global.store %arg0, @unfusableDivergent0 : index util.global.store %arg0, @unfusableDivergent1 : index return } -builtin.func @fn_b(%arg0: index) { +func.func @fn_b(%arg0: index) { util.global.store %arg0, @unfusableDivergent0 : index return }
diff --git a/iree/compiler/Dialect/Util/Transforms/test/hoist_into_globals.mlir b/iree/compiler/Dialect/Util/Transforms/test/hoist_into_globals.mlir index 002c0b6..567890a 100644 --- a/iree/compiler/Dialect/Util/Transforms/test/hoist_into_globals.mlir +++ b/iree/compiler/Dialect/Util/Transforms/test/hoist_into_globals.mlir
@@ -4,7 +4,7 @@ module @hoist_simple_const_expr { // CHECK: util.global private @[[HOISTED_SYM:.*]] : i32 // CHECK: func @main - builtin.func @main() -> (i32) { + func.func @main() -> (i32) { %0 = arith.constant 0 : i32 %1 = arith.constant 1 : i32 // CHECK-NOT: arith.constant @@ -33,7 +33,7 @@ // CHECK: return %[[VAL]] // CHECK-NOT: util.initializer module @do_not_hoist_variable_op { - builtin.func @main() -> (i32) { + func.func @main() -> (i32) { %0 = arith.constant 0 : i32 %1 = arith.constant 1 : i32 %2 = "iree_unregistered.var_expr"(%0, %1) : (i32, i32) -> i32 @@ -46,7 +46,7 @@ // CHECK-NOT: util.global // CHECK-NOT: util.initializer module @do_not_hoist_variable_operands { - builtin.func @main(%arg0 : i32) -> (i32) { + func.func @main(%arg0 : i32) -> (i32) { %0 = arith.constant 0 : i32 %2 = "iree_unregistered.const_expr"(%0, %arg0) : (i32, i32) -> i32 return %2 : i32 @@ -58,7 +58,7 @@ // CHECK-NOT: util.global // CHECK-NOT: util.initializer module @do_not_hoist_sub_byte_aligned_scalar_leaf { - builtin.func @main() -> (i32) { + func.func @main() -> (i32) { %0 = arith.constant 1 : i1 %2 = "iree_unregistered.var_expr"(%0) : (i1) -> i32 return %2 : i32 @@ -70,7 +70,7 @@ // CHECK-NOT: util.global // CHECK-NOT: util.initializer module @do_not_hoist_sub_byte_aligned_tensor_leaf { - builtin.func @main() -> (i32) { + func.func @main() -> (i32) { %0 = arith.constant dense<true> : tensor<i1> %2 = "iree_unregistered.var_expr"(%0) : (tensor<i1>) -> i32 return %2 : i32 @@ -83,7 +83,7 @@ // Can hoist a const-expr tree that transitively includes sub-byte aligned // values. module @hoist_sub_byte_aligned_scalar_transitive { - builtin.func @main() -> (i32) { + func.func @main() -> (i32) { %0 = arith.constant 1 : i1 %2 = "iree_unregistered.const_expr"(%0) : (i1) -> i32 return %2 : i32 @@ -96,7 +96,7 @@ // Can hoist a const-expr tree that transitively includes sub-byte aligned // values. module @hoist_sub_byte_aligned_tensor_transitive { - builtin.func @main() -> (i32) { + func.func @main() -> (i32) { %0 = arith.constant dense<true> : tensor<i1> %2 = "iree_unregistered.const_expr"(%0) : (tensor<i1>) -> i32 return %2 : i32 @@ -113,7 +113,7 @@ util.global private @latent_global : i32 // CHECK: func @main - builtin.func @main() -> (i32, i32, i32) { + func.func @main() -> (i32, i32, i32) { // CHECK-DAG: %[[LOAD_HOISTED_0:.*]] = util.global.load @[[HOISTED_0]] : i32 // CHECK-DAG: %[[LOAD_HOISTED_1:.*]] = util.global.load @[[HOISTED_1]] : i32 // CHECK-DAG: %[[RESULT:.*]] = "iree_unregistered.var_expr"(%[[LOAD_HOISTED_1]]) @@ -150,7 +150,7 @@ module @hoist_non_leaf_const_expr { // CHECK: util.global private @[[HOISTED:.*]] : i32 // CHECK: func @main - builtin.func @main() -> (i32) { + func.func @main() -> (i32) { // CHECK: %[[LOAD_HOISTED:.*]] = util.global.load @[[HOISTED]] : i32 // CHECK: %[[RESULT:.*]] = "iree_unregistered.non_leaf_const_expr"(%hoisted) // CHECK: return %[[RESULT]] @@ -176,7 +176,7 @@ module @hoist_implicit_capture { // CHECK: util.global private @[[HOISTED_SYM:.*]] : i32 // CHECK: func @main - builtin.func @main() -> (i32) { + func.func @main() -> (i32) { %0 = arith.constant 0 : i32 %1 = arith.constant 1 : i32 // CHECK-NOT: arith.constant
diff --git a/iree/compiler/Dialect/Util/Transforms/test/hoist_into_globals_linalg.mlir b/iree/compiler/Dialect/Util/Transforms/test/hoist_into_globals_linalg.mlir index 871ac46..ef5ae65 100644 --- a/iree/compiler/Dialect/Util/Transforms/test/hoist_into_globals_linalg.mlir +++ b/iree/compiler/Dialect/Util/Transforms/test/hoist_into_globals_linalg.mlir
@@ -7,7 +7,7 @@ module @compute_hoisted { // CHECK: util.global private @[[HOISTED:.*]] : tensor<5x6xf32> // CHECK: func @main - builtin.func @main() -> (tensor<5x6xf32>) { + func.func @main() -> (tensor<5x6xf32>) { %cst_0 = arith.constant dense<1.270000e+02> : tensor<f32> // A non-leaf broadcast. @@ -42,7 +42,7 @@ module @broadcast_treated_as_leaf { // CHECK-NOT: util.global // CHECK: func @main - builtin.func @main() -> (tensor<5x6xf32>) { + func.func @main() -> (tensor<5x6xf32>) { %cst_0 = arith.constant dense<1.270000e+02> : tensor<f32> // CHECK: linalg.init_tensor %0 = linalg.init_tensor [5, 6] : tensor<5x6xf32>
diff --git a/iree/compiler/Dialect/Util/Transforms/test/simplify_global_accesses.mlir b/iree/compiler/Dialect/Util/Transforms/test/simplify_global_accesses.mlir index bca8c4a..59bed47 100644 --- a/iree/compiler/Dialect/Util/Transforms/test/simplify_global_accesses.mlir +++ b/iree/compiler/Dialect/Util/Transforms/test/simplify_global_accesses.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='builtin.func(iree-util-simplify-global-accesses)' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='func.func(iree-util-simplify-global-accesses)' %s | FileCheck %s util.global private @varA = dense<1> : tensor<2xi32> util.global private @varB = dense<3> : tensor<2x4xi32>
diff --git a/iree/compiler/Dialect/Util/Transforms/test/strip_debug_ops.mlir b/iree/compiler/Dialect/Util/Transforms/test/strip_debug_ops.mlir index d4c518c..e6a1245 100644 --- a/iree/compiler/Dialect/Util/Transforms/test/strip_debug_ops.mlir +++ b/iree/compiler/Dialect/Util/Transforms/test/strip_debug_ops.mlir
@@ -1,4 +1,4 @@ -// RUN: iree-opt -split-input-file -pass-pipeline='builtin.func(iree-util-strip-debug-ops)' %s | FileCheck %s +// RUN: iree-opt -split-input-file -pass-pipeline='func.func(iree-util-strip-debug-ops)' %s | FileCheck %s // CHECK-LABEL: @stripAssert func @stripAssert(%cond: i1) {
diff --git a/iree/compiler/Dialect/VM/Conversion/ImportUtils.h b/iree/compiler/Dialect/VM/Conversion/ImportUtils.h index bf2e057..cf2440b 100644 --- a/iree/compiler/Dialect/VM/Conversion/ImportUtils.h +++ b/iree/compiler/Dialect/VM/Conversion/ImportUtils.h
@@ -57,7 +57,7 @@ state.addAttributes(llvm::to_vector<4>(operation->getDialectAttrs())); state.addAttribute("callee", SymbolRefAttr::get(importOp)); - auto importType = importOp.getType(); + auto importType = importOp.getFunctionType(); for (auto resultType : operation->getResultTypes()) { if (failed(typeConverter.convertType(resultType, state.types))) { return None;
diff --git a/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVM.cpp b/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVM.cpp index 8f07600..aa2c873 100644 --- a/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVM.cpp +++ b/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVM.cpp
@@ -69,7 +69,7 @@ LogicalResult matchAndRewrite( FuncOp srcOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - FunctionType srcFuncType = srcOp.getType(); + FunctionType srcFuncType = srcOp.getFunctionType(); TypeConverter::SignatureConversion signatureConversion( srcOp.getNumArguments());
diff --git a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp index f3cc426..bab0054 100644 --- a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp +++ b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp
@@ -135,7 +135,7 @@ auto moduleOp = funcOp.getOperation()->getParentOfType<IREE::VM::ModuleOp>(); - FunctionType funcType = funcOp.getType(); + FunctionType funcType = funcOp.getFunctionType(); std::string name = std::string(moduleOp.getName()) + "_" + std::string(funcOp.getName()); std::string moduleTypeName = (moduleOp.getName() + "_t").str(); @@ -172,7 +172,7 @@ auto newFuncType = mlir::FunctionType::get( ctx, {inputTypes}, {emitc::OpaqueType::get(ctx, "iree_status_t")}); - auto newFuncOp = builder.create<mlir::FuncOp>(loc, name, newFuncType); + auto newFuncOp = builder.create<mlir::func::FuncOp>(loc, name, newFuncType); newFuncOp.getOperation()->setAttr("emitc.static", UnitAttr::get(ctx)); @@ -214,7 +214,7 @@ vmAnalysis = typeConverter.lookupAnalysis(newFuncOp); if (failed(vmAnalysis)) { return funcOp.emitError() - << "newly created mlir::FuncOp not found in cache."; + << "newly created mlir::func::FuncOp not found in cache."; } // Add constant ops for local refs @@ -458,7 +458,8 @@ } /// Releases refs which are local to the function as well as ref arguments. -void releaseRefs(OpBuilder &builder, Location location, mlir::FuncOp funcOp, +void releaseRefs(OpBuilder &builder, Location location, + mlir::func::FuncOp funcOp, IREE::VM::EmitCTypeConverter &typeConverter) { auto ctx = builder.getContext(); @@ -567,7 +568,7 @@ auto blockBuilder = [&builder, &location, &typeConverter](emitc::CallOp &callOp) { Block *block = builder.getBlock(); - mlir::FuncOp funcOp = cast<mlir::FuncOp>(block->getParentOp()); + mlir::func::FuncOp funcOp = cast<mlir::func::FuncOp>(block->getParentOp()); releaseRefs(builder, location, funcOp, typeConverter); @@ -589,7 +590,7 @@ auto ctx = builder.getContext(); Block *block = builder.getBlock(); - mlir::FuncOp funcOp = cast<mlir::FuncOp>(block->getParentOp()); + mlir::func::FuncOp funcOp = cast<mlir::func::FuncOp>(block->getParentOp()); releaseRefs(builder, location, funcOp, typeConverter); @@ -615,7 +616,7 @@ /// value, i.e. a truthy value branches to the continuation block when /// `negateCondition` is false. mlir::func::CallOp failableCall( - OpBuilder &builder, Location location, mlir::FuncOp &callee, + OpBuilder &builder, Location location, mlir::func::FuncOp &callee, ArrayRef<Value> operands, const std::function<void(mlir::func::CallOp &)> &failureBlockBuilder, bool negateCondition = false) { @@ -665,12 +666,13 @@ } mlir::func::CallOp returnIfError(OpBuilder &builder, Location location, - mlir::FuncOp &callee, ArrayRef<Value> operands, + mlir::func::FuncOp &callee, + ArrayRef<Value> operands, IREE::VM::EmitCTypeConverter &typeConverter) { auto blockBuilder = [&builder, &location, &typeConverter](mlir::func::CallOp &callOp) { Block *block = builder.getBlock(); - mlir::FuncOp funcOp = cast<mlir::FuncOp>(block->getParentOp()); + mlir::func::FuncOp funcOp = cast<mlir::func::FuncOp>(block->getParentOp()); releaseRefs(builder, location, funcOp, typeConverter); @@ -699,8 +701,8 @@ ctx, {emitc::PointerType::get(emitc::OpaqueType::get(ctx, "void"))}, {}); - auto funcOp = - builder.create<mlir::FuncOp>(loc, moduleName + "_destroy", funcType); + auto funcOp = builder.create<mlir::func::FuncOp>( + loc, moduleName + "_destroy", funcType); typeConverter.analysisCache.insert( std::make_pair(funcOp.getOperation(), VMAnalysis())); @@ -760,8 +762,8 @@ emitc::OpaqueType::get(ctx, "iree_vm_module_state_t")))}, {emitc::OpaqueType::get(ctx, "iree_status_t")}); - auto funcOp = builder.create<mlir::FuncOp>(loc, moduleName + "_alloc_state", - funcType); + auto funcOp = builder.create<mlir::func::FuncOp>( + loc, moduleName + "_alloc_state", funcType); typeConverter.analysisCache.insert( std::make_pair(funcOp.getOperation(), VMAnalysis())); @@ -1012,8 +1014,8 @@ emitc::OpaqueType::get(ctx, "iree_vm_module_state_t"))}, {}); - auto funcOp = - builder.create<mlir::FuncOp>(loc, moduleName + "_free_state", funcType); + auto funcOp = builder.create<mlir::func::FuncOp>( + loc, moduleName + "_free_state", funcType); typeConverter.analysisCache.insert( std::make_pair(funcOp.getOperation(), VMAnalysis())); @@ -1125,7 +1127,7 @@ }, {emitc::OpaqueType::get(ctx, "iree_status_t")}); - auto funcOp = builder.create<mlir::FuncOp>( + auto funcOp = builder.create<mlir::func::FuncOp>( loc, moduleName + "_resolve_import", funcType); typeConverter.analysisCache.insert( @@ -1207,8 +1209,8 @@ emitc::OpaqueType::get(ctx, "iree_vm_module_t")))}, {emitc::OpaqueType::get(ctx, "iree_status_t")}); - auto funcOp = - builder.create<mlir::FuncOp>(loc, moduleName + "_create", funcType); + auto funcOp = builder.create<mlir::func::FuncOp>( + loc, moduleName + "_create", funcType); typeConverter.analysisCache.insert( std::make_pair(funcOp.getOperation(), VMAnalysis())); @@ -1464,16 +1466,16 @@ StringRef funcName; }; -class FuncOpConversion : public OpConversionPattern<mlir::FuncOp> { +class FuncOpConversion : public OpConversionPattern<mlir::func::FuncOp> { public: - using OpConversionPattern<mlir::FuncOp>::OpConversionPattern; + using OpConversionPattern<mlir::func::FuncOp>::OpConversionPattern; private: LogicalResult matchAndRewrite( - mlir::FuncOp funcOp, OpAdaptor adaptor, + mlir::func::FuncOp funcOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { TypeConverter::SignatureConversion signatureConverter( - funcOp.getType().getNumInputs()); + funcOp.getFunctionType().getNumInputs()); TypeConverter typeConverter; for (const auto &arg : llvm::enumerate(funcOp.getArguments())) { Type convertedType = @@ -1487,7 +1489,7 @@ rewriter.updateRootInPlace(funcOp, [&] { funcOp.setType( rewriter.getFunctionType(signatureConverter.getConvertedTypes(), - funcOp.getType().getResults())); + funcOp.getFunctionType().getResults())); }); return success(); } @@ -1520,8 +1522,8 @@ rewriter.startRootUpdate(exportOp.getOperation()); - mlir::FuncOp funcOp = - lookupSymbolRef<mlir::FuncOp>(exportOp.getOperation(), "function_ref"); + mlir::func::FuncOp funcOp = lookupSymbolRef<mlir::func::FuncOp>( + exportOp.getOperation(), "function_ref"); auto vmAnalysis = typeConverter->lookupAnalysis(funcOp); if (failed(vmAnalysis)) { @@ -1554,7 +1556,7 @@ ctx, {inputTypes}, {emitc::OpaqueType::get(ctx, "iree_status_t")}); auto newFuncOp = - rewriter.create<mlir::FuncOp>(loc, newFuncName, newFuncType); + rewriter.create<mlir::func::FuncOp>(loc, newFuncName, newFuncType); VMAnalysis newVmAnalysis; newVmAnalysis.numRefArguments = numRefArgs; @@ -1659,7 +1661,7 @@ FailureOr<std::pair<Value, Value>> castModuleAndStateStructs( ConversionPatternRewriter &rewriter, IREE::VM::ExportOp &exportOp, - mlir::FuncOp &newFuncOp) const { + mlir::func::FuncOp &newFuncOp) const { auto ctx = exportOp.getContext(); auto loc = exportOp.getLoc(); @@ -1705,14 +1707,14 @@ FailureOr<std::pair<GeneratedStruct, GeneratedStruct>> typedefArgumentAndResultStructs(ConversionPatternRewriter &rewriter, IREE::VM::ExportOp &exportOp, - mlir::FuncOp &newFuncOp) const { + mlir::func::FuncOp &newFuncOp) const { auto loc = exportOp.getLoc(); IREE::VM::EmitCTypeConverter *typeConverter = this->template getTypeConverter<IREE::VM::EmitCTypeConverter>(); - mlir::FuncOp funcOp = - lookupSymbolRef<mlir::FuncOp>(exportOp.getOperation(), "function_ref"); + mlir::func::FuncOp funcOp = lookupSymbolRef<mlir::func::FuncOp>( + exportOp.getOperation(), "function_ref"); auto vmAnalysis = typeConverter->lookupAnalysis(funcOp); if (failed(vmAnalysis)) { @@ -1802,7 +1804,7 @@ void castArgumentAndResultStructs(ConversionPatternRewriter &rewriter, IREE::VM::ExportOp &exportOp, - mlir::FuncOp &newFuncOp, + mlir::func::FuncOp &newFuncOp, GeneratedStruct &argumentStruct, GeneratedStruct &resultStruct) const { auto ctx = exportOp.getContext(); @@ -1909,8 +1911,8 @@ IREE::VM::EmitCTypeConverter *typeConverter = this->template getTypeConverter<IREE::VM::EmitCTypeConverter>(); - mlir::FuncOp funcOp = - lookupSymbolRef<mlir::FuncOp>(exportOp.getOperation(), "function_ref"); + mlir::func::FuncOp funcOp = lookupSymbolRef<mlir::func::FuncOp>( + exportOp.getOperation(), "function_ref"); auto vmAnalysis = typeConverter->lookupAnalysis(funcOp); if (failed(vmAnalysis)) { @@ -1972,8 +1974,8 @@ IREE::VM::EmitCTypeConverter *typeConverter = this->template getTypeConverter<IREE::VM::EmitCTypeConverter>(); - mlir::FuncOp funcOp = - lookupSymbolRef<mlir::FuncOp>(exportOp.getOperation(), "function_ref"); + mlir::func::FuncOp funcOp = lookupSymbolRef<mlir::func::FuncOp>( + exportOp.getOperation(), "function_ref"); auto vmAnalysis = typeConverter->lookupAnalysis(funcOp); if (failed(vmAnalysis)) { @@ -2038,13 +2040,13 @@ } if (importOp.isVariadic()) { - if (failed(createVariadicImportShims(importOp.getType(), importOp, + if (failed(createVariadicImportShims(importOp.getFunctionType(), importOp, rewriter))) { return failure(); } } else { - if (failed( - createImportShim(importOp.getType(), importOp, -1, rewriter))) { + if (failed(createImportShim(importOp.getFunctionType(), importOp, -1, + rewriter))) { return failure(); } } @@ -2101,8 +2103,8 @@ << "Failed to build function type for wrapper"; } - auto newFuncOp = rewriter.create<mlir::FuncOp>(loc, newFuncName.getValue(), - newFuncType.getValue()); + auto newFuncOp = rewriter.create<mlir::func::FuncOp>( + loc, newFuncName.getValue(), newFuncType.getValue()); getTypeConverter<IREE::VM::EmitCTypeConverter>()->analysisCache.insert( std::make_pair(newFuncOp.getOperation(), VMAnalysis{})); @@ -2115,7 +2117,7 @@ Block *block = rewriter.createBlock(&newFuncOp.getBody(), newFuncOp.getBody().end()); - for (Type type : newFuncOp.getType().getInputs()) { + for (Type type : newFuncOp.getFunctionType().getInputs()) { block->addArgument(type, loc); } @@ -2395,7 +2397,7 @@ } LogicalResult packArgumentBuffer(ArrayRef<Type> inputTypes, - mlir::FuncOp &funcOp, Value call, + mlir::func::FuncOp &funcOp, Value call, ConversionPatternRewriter &rewriter, Location loc) const { auto ctx = rewriter.getContext(); @@ -2520,7 +2522,7 @@ } LogicalResult unpackResultBuffer(ArrayRef<Type> resultTypes, - mlir::FuncOp &funcOp, Value call, + mlir::func::FuncOp &funcOp, Value call, ConversionPatternRewriter &rewriter, Location loc) const { auto ctx = rewriter.getContext(); @@ -2785,8 +2787,8 @@ LogicalResult matchAndRewrite( CallOpTy op, Adaptor adaptor, ConversionPatternRewriter &rewriter) const override { - mlir::FuncOp funcOp = - lookupSymbolRef<mlir::FuncOp>(op.getOperation(), "callee"); + mlir::func::FuncOp funcOp = + lookupSymbolRef<mlir::func::FuncOp>(op.getOperation(), "callee"); IREE::VM::ImportOp importOp = lookupSymbolRef<IREE::VM::ImportOp>(op.getOperation(), "callee"); @@ -2806,13 +2808,13 @@ LogicalResult rewriteInternalCall(Operation *op, Adaptor adaptor, ConversionPatternRewriter &rewriter, - mlir::FuncOp funcOp) const { + mlir::func::FuncOp funcOp) const { auto loc = op->getLoc(); SmallVector<Value, 4> updatedOperands; SmallVector<Value, 4> resultOperands; - auto parentFuncOp = op->getParentOfType<mlir::FuncOp>(); + auto parentFuncOp = op->getParentOfType<mlir::func::FuncOp>(); BlockArgument stackArg = parentFuncOp.getArgument(0); BlockArgument moduleArg = parentFuncOp.getArgument(1); @@ -2854,7 +2856,7 @@ int importOrdinal = importOp.ordinal().getValue().getZExtValue(); - auto funcOp = op->getParentOfType<mlir::FuncOp>(); + auto funcOp = op->getParentOfType<mlir::func::FuncOp>(); BlockArgument stackArg = funcOp.getArgument(0); BlockArgument stateArg = funcOp.getArgument(2); @@ -2913,7 +2915,8 @@ if (!funcName.hasValue()) return op->emitError() << "Couldn't build name to imported function"; - auto callee = moduleOp.lookupSymbol<mlir::FuncOp>(funcName.getValue()); + auto callee = + moduleOp.lookupSymbol<mlir::func::FuncOp>(funcName.getValue()); if (callee == nullptr) { return op->emitError() << "Couldn't find function with name `" << funcName.getValue() << "`"; @@ -3086,7 +3089,7 @@ auto loc = cmpOp.getLoc(); auto funcOp = - cmpOp.getOperation()->template getParentOfType<mlir::FuncOp>(); + cmpOp.getOperation()->template getParentOfType<mlir::func::FuncOp>(); IREE::VM::EmitCTypeConverter *typeConverter = this->template getTypeConverter<IREE::VM::EmitCTypeConverter>(); @@ -3161,7 +3164,7 @@ auto ctx = cmpOp.getContext(); auto loc = cmpOp.getLoc(); - auto funcOp = cmpOp.getOperation()->getParentOfType<mlir::FuncOp>(); + auto funcOp = cmpOp.getOperation()->getParentOfType<mlir::func::FuncOp>(); IREE::VM::EmitCTypeConverter *typeConverter = this->template getTypeConverter<IREE::VM::EmitCTypeConverter>(); @@ -3296,7 +3299,7 @@ } auto funcOp = - constRefRodataOp.getOperation()->getParentOfType<mlir::FuncOp>(); + constRefRodataOp.getOperation()->getParentOfType<mlir::func::FuncOp>(); BlockArgument stateArg = funcOp.getArgument(2); auto rodataBuffersPtr = rewriter.create<emitc::CallOp>( @@ -3393,7 +3396,7 @@ return success(); } - auto funcOp = op.getOperation()->getParentOfType<mlir::FuncOp>(); + auto funcOp = op.getOperation()->getParentOfType<mlir::func::FuncOp>(); IREE::VM::EmitCTypeConverter *typeConverter = this->template getTypeConverter<IREE::VM::EmitCTypeConverter>(); @@ -3522,7 +3525,7 @@ return success(); } - auto funcOp = op.getOperation()->getParentOfType<mlir::FuncOp>(); + auto funcOp = op.getOperation()->getParentOfType<mlir::func::FuncOp>(); IREE::VM::EmitCTypeConverter *typeConverter = getTypeConverter<IREE::VM::EmitCTypeConverter>(); @@ -3632,7 +3635,7 @@ auto ctx = op.getContext(); auto loc = op.getLoc(); - auto funcOp = op.getOperation()->getParentOfType<mlir::FuncOp>(); + auto funcOp = op.getOperation()->getParentOfType<mlir::func::FuncOp>(); IREE::VM::EmitCTypeConverter *typeConverter = getTypeConverter<IREE::VM::EmitCTypeConverter>(); @@ -3755,7 +3758,7 @@ passthroughBlock = rewriter.createBlock(parentRegion, parentRegion->end()); - auto funcOp = op.getOperation()->getParentOfType<mlir::FuncOp>(); + auto funcOp = op.getOperation()->getParentOfType<mlir::func::FuncOp>(); IREE::VM::EmitCTypeConverter *typeConverter = getTypeConverter<IREE::VM::EmitCTypeConverter>(); @@ -3776,7 +3779,7 @@ OpBuilder::InsertionGuard guard(rewriter); failureBlock = rewriter.createBlock(parentRegion, parentRegion->end()); - auto funcOp = op.getOperation()->getParentOfType<mlir::FuncOp>(); + auto funcOp = op.getOperation()->getParentOfType<mlir::func::FuncOp>(); IREE::VM::EmitCTypeConverter *typeConverter = getTypeConverter<IREE::VM::EmitCTypeConverter>(); @@ -3889,7 +3892,7 @@ } auto funcOp = - loadOp.getOperation()->template getParentOfType<mlir::FuncOp>(); + loadOp.getOperation()->template getParentOfType<mlir::func::FuncOp>(); BlockArgument stateArg = funcOp.getArgument(2); auto rwDataPtr = rewriter.create<emitc::CallOp>( @@ -3955,7 +3958,7 @@ auto globalOrdinal = globalOp.ordinal().getValue().getZExtValue(); - auto funcOp = op->getParentOfType<mlir::FuncOp>(); + auto funcOp = op->getParentOfType<mlir::func::FuncOp>(); IREE::VM::EmitCTypeConverter *typeConverter = this->template getTypeConverter<IREE::VM::EmitCTypeConverter>(); @@ -4069,7 +4072,7 @@ } auto funcOp = - storeOp.getOperation()->template getParentOfType<mlir::FuncOp>(); + storeOp.getOperation()->template getParentOfType<mlir::func::FuncOp>(); BlockArgument stateArg = funcOp.getArgument(2); auto rwDataPtr = rewriter.create<emitc::CallOp>( @@ -4240,7 +4243,7 @@ /*applicableOperator=*/StringAttr::get(ctx, "&"), /*operand=*/listOp.getResult()); - auto funcOp = allocOp.getOperation()->getParentOfType<mlir::FuncOp>(); + auto funcOp = allocOp.getOperation()->getParentOfType<mlir::func::FuncOp>(); IREE::VM::EmitCTypeConverter *typeConverter = getTypeConverter<IREE::VM::EmitCTypeConverter>(); @@ -4671,7 +4674,7 @@ /*operands=*/ArrayRef<Value>{refOp.getResult()}, /*typeConverter=*/*typeConverter); - auto funcOp = setOp.getOperation()->getParentOfType<mlir::FuncOp>(); + auto funcOp = setOp.getOperation()->getParentOfType<mlir::func::FuncOp>(); auto vmAnalysis = typeConverter->lookupAnalysis(funcOp); if (failed(vmAnalysis)) { @@ -5096,9 +5099,10 @@ mlir::func::FuncDialect, mlir::arith::ArithmeticDialect, mlir::math::MathDialect>(); - target.addDynamicallyLegalOp<mlir::FuncOp>([&](mlir::FuncOp op) { - return typeConverter.isSignatureLegal(op.getType()); - }); + target.addDynamicallyLegalOp<mlir::func::FuncOp>( + [&](mlir::func::FuncOp op) { + return typeConverter.isSignatureLegal(op.getFunctionType()); + }); // Structural ops target.addLegalOp<IREE::VM::ModuleOp>();
diff --git a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/EmitCTypeConverter.cpp b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/EmitCTypeConverter.cpp index 7d2dfab..516b532 100644 --- a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/EmitCTypeConverter.cpp +++ b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/EmitCTypeConverter.cpp
@@ -80,12 +80,12 @@ Optional<Value> EmitCTypeConverter::materializeRef(Value ref) { assert(ref.getType().isa<IREE::VM::RefType>()); - mlir::FuncOp funcOp; + mlir::func::FuncOp funcOp; if (auto definingOp = ref.getDefiningOp()) { - funcOp = definingOp->getParentOfType<mlir::FuncOp>(); + funcOp = definingOp->getParentOfType<mlir::func::FuncOp>(); } else { Operation *op = ref.cast<BlockArgument>().getOwner()->getParentOp(); - funcOp = cast<mlir::FuncOp>(op); + funcOp = cast<mlir::func::FuncOp>(op); } auto vmAnalysis = lookupAnalysis(funcOp);
diff --git a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/EmitCTypeConverter.h b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/EmitCTypeConverter.h index 59bfec0..0856a5b 100644 --- a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/EmitCTypeConverter.h +++ b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/EmitCTypeConverter.h
@@ -10,6 +10,7 @@ #include "iree/compiler/Dialect/VM/Conversion/VMToEmitC/VMAnalysis.h" #include "iree/compiler/Dialect/VM/IR/VMTypes.h" #include "mlir/Dialect/EmitC/IR/EmitC.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Transforms/DialectConversion.h" namespace mlir { @@ -21,7 +22,7 @@ public: EmitCTypeConverter(); FailureOr<std::reference_wrapper<VMAnalysis>> lookupAnalysis( - mlir::FuncOp &funcOp) { + mlir::func::FuncOp &funcOp) { return lookupAnalysis(funcOp.getOperation()); } FailureOr<std::reference_wrapper<VMAnalysis>> lookupAnalysis(
diff --git a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/VMAnalysis.h b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/VMAnalysis.h index 628a7aa..e8f93e4 100644 --- a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/VMAnalysis.h +++ b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/VMAnalysis.h
@@ -22,7 +22,7 @@ Operation *op = funcOp.getOperation(); registerAllocation = RegisterAllocation(op); valueLiveness = ValueLiveness(op); - originalFunctionType = funcOp.getType(); + originalFunctionType = funcOp.getFunctionType(); } VMAnalysis(VMAnalysis &&) = default;
diff --git a/iree/compiler/Dialect/VM/IR/VMOps.cpp b/iree/compiler/Dialect/VM/IR/VMOps.cpp index 9dfa11c..d9a2064 100644 --- a/iree/compiler/Dialect/VM/IR/VMOps.cpp +++ b/iree/compiler/Dialect/VM/IR/VMOps.cpp
@@ -86,9 +86,7 @@ void FuncOp::print(OpAsmPrinter &p) { Operation *op = getOperation(); - FunctionType fnType = getType(); - function_interface_impl::printFunctionOp( - p, op, fnType.getInputs(), /*isVariadic=*/false, fnType.getResults()); + function_interface_impl::printFunctionOp(p, op, /*isVariadic=*/false); } void FuncOp::build(OpBuilder &builder, OperationState &result, StringRef name, @@ -97,7 +95,7 @@ result.addRegion(); result.addAttribute(SymbolTable::getSymbolAttrName(), builder.getStringAttr(name)); - result.addAttribute("type", TypeAttr::get(type)); + result.addAttribute("function_type", TypeAttr::get(type)); result.attributes.append(attrs.begin(), attrs.end()); if (argAttrs.empty()) { return; @@ -113,13 +111,14 @@ assert(empty() && "function already has an entry block"); auto *entry = new Block(); push_back(entry); - SmallVector<Location> locs(getType().getNumInputs(), getLoc()); - entry->addArguments(getType().getInputs(), locs); + SmallVector<Location> locs(getFunctionType().getNumInputs(), getLoc()); + entry->addArguments(getFunctionType().getInputs(), locs); return entry; } LogicalResult FuncOp::verifyType() { - auto type = getTypeAttr().getValue(); + auto type = + getOperation()->getAttrOfType<TypeAttr>(getTypeAttrName()).getValue(); if (!type.isa<FunctionType>()) return emitOpError("requires '" + getTypeAttrName() + "' attribute of function type"); @@ -268,7 +267,7 @@ if (auto name = getArgAttrOfType<StringAttr>(i, "vm.name")) { p << '%' << name.getValue() << " : "; } - p.printType(getType().getInput(i)); + p.printType(getFunctionType().getInput(i)); if (getArgAttrOfType<UnitAttr>(i, "vm.variadic")) { p << " ..."; } @@ -279,9 +278,9 @@ p << ")"; if (getResultTypes().size() == 1) { p << " -> "; - p.printType(getType().getResult(0)); + p.printType(getFunctionType().getResult(0)); } else if (getResultTypes().size() > 1) { - p << " -> (" << getType().getResults() << ")"; + p << " -> (" << getFunctionType().getResults() << ")"; } mlir::function_interface_impl::printFunctionAttributes( p, op, getArgumentTypes().size(), getResultTypes().size(), @@ -296,7 +295,7 @@ ArrayRef<DictionaryAttr> argAttrs) { result.addAttribute(SymbolTable::getSymbolAttrName(), builder.getStringAttr(name)); - result.addAttribute("type", TypeAttr::get(type)); + result.addAttribute("function_type", TypeAttr::get(type)); result.attributes.append(attrs.begin(), attrs.end()); if (argAttrs.empty()) { return; @@ -309,7 +308,8 @@ } LogicalResult ImportOp::verifyType() { - auto type = getTypeAttr().getValue(); + auto type = + getOperation()->getAttrOfType<TypeAttr>(getTypeAttrName()).getValue(); if (!type.isa<FunctionType>()) return emitOpError("requires '" + getTypeAttrName() + "' attribute of function type"); @@ -318,15 +318,15 @@ void InitializerOp::build(OpBuilder &builder, OperationState &result, ArrayRef<NamedAttribute> attrs) { - result.addAttribute( - "type", TypeAttr::get(FunctionType::get(builder.getContext(), {}, {}))); + result.addAttribute("function_type", TypeAttr::get(FunctionType::get( + builder.getContext(), {}, {}))); result.addRegion(); result.attributes.append(attrs.begin(), attrs.end()); } ParseResult InitializerOp::parse(OpAsmParser &parser, OperationState &result) { - result.addAttribute( - "type", TypeAttr::get(FunctionType::get(result.getContext(), {}, {}))); + result.addAttribute("function_type", TypeAttr::get(FunctionType::get( + result.getContext(), {}, {}))); if (parser.parseOptionalAttrDictWithKeyword(result.attributes)) { return failure(); } @@ -339,7 +339,8 @@ void InitializerOp::print(OpAsmPrinter &p) { Operation *op = getOperation(); - p.printOptionalAttrDictWithKeyword(op->getAttrs(), /*elidedAttrs=*/{"type"}); + p.printOptionalAttrDictWithKeyword(op->getAttrs(), + /*elidedAttrs=*/{"function_type"}); p << " "; p.printRegion(body()); } @@ -378,13 +379,14 @@ return op->emitOpError() << "initializer function " << initializerAttr << " not found"; } - if (initializer.getType().getNumInputs() != 0 || - initializer.getType().getNumResults() != 1 || - initializer.getType().getResult(0) != globalType.getValue()) { + if (initializer.getFunctionType().getNumInputs() != 0 || + initializer.getFunctionType().getNumResults() != 1 || + initializer.getFunctionType().getResult(0) != globalType.getValue()) { return op->emitOpError() << "initializer type mismatch; global " << globalName << " is " << globalType << " but initializer function " - << initializer.getName() << " is " << initializer.getType(); + << initializer.getName() << " is " + << initializer.getFunctionType(); } } else if (initialValueAttr) { // Ensure the value is something we can convert to a const.
diff --git a/iree/compiler/Dialect/VM/IR/VMOps.td b/iree/compiler/Dialect/VM/IR/VMOps.td index 9c06818..2d21ab8 100644 --- a/iree/compiler/Dialect/VM/IR/VMOps.td +++ b/iree/compiler/Dialect/VM/IR/VMOps.td
@@ -114,17 +114,17 @@ } /// Returns the type of this function. - FunctionType getType() { + FunctionType getFunctionType() { return getOperation()->getAttrOfType<TypeAttr>(getTypeAttrName()) .getValue() .cast<FunctionType>(); } /// Returns the argument types of this function. - ArrayRef<Type> getArgumentTypes() { return getType().getInputs(); } + ArrayRef<Type> getArgumentTypes() { return getFunctionType().getInputs(); } /// Returns the result types of this function. - ArrayRef<Type> getResultTypes() { return getType().getResults(); } + ArrayRef<Type> getResultTypes() { return getFunctionType().getResults(); } /// Hook for OpTrait::FunctionLike, called after verifying that the 'type' /// attribute is present. This can check for preconditions of the @@ -134,7 +134,7 @@ Region *getCallableRegion() { return &body(); } ArrayRef<Type> getCallableResults() { assert(!isExternal() && "invalid callable"); - return getType().getResults(); + return getFunctionType().getResults(); } /// Adds or overrides a reflection attribute. @@ -195,7 +195,7 @@ let extraClassDeclaration = [{ /// Returns the type of this function. - FunctionType getType() { + FunctionType getFunctionType() { return getOperation()->getAttrOfType<TypeAttr>(getTypeAttrName()) .getValue() .cast<FunctionType>(); @@ -220,10 +220,10 @@ } /// Returns the argument types of this function. - ArrayRef<Type> getArgumentTypes() { return getType().getInputs(); } + ArrayRef<Type> getArgumentTypes() { return getFunctionType().getInputs(); } /// Returns the result types of this function. - ArrayRef<Type> getResultTypes() { return getType().getResults(); } + ArrayRef<Type> getResultTypes() { return getFunctionType().getResults(); } /// Hook for OpTrait::FunctionLike, called after verifying that the 'type' @@ -233,7 +233,7 @@ Region *getCallableRegion() { return nullptr; } ArrayRef<Type> getCallableResults() { - return getType().getResults(); + return getFunctionType().getResults(); } }]; } @@ -252,7 +252,7 @@ }]; let arguments = (ins - TypeAttr:$type + TypeAttr:$function_type ); let regions = (region AnyRegion:$body); @@ -270,8 +270,10 @@ Block *addEntryBlock(); Block *addBlock(); - FunctionType getType() { - return getTypeAttr().getValue().cast<FunctionType>(); + FunctionType getFunctionType() { + return getOperation()->getAttrOfType<TypeAttr>(getTypeAttrName()) + .getValue() + .cast<FunctionType>(); } /// Returns the argument types of this function. @@ -3530,7 +3532,7 @@ [{ $_state.addOperands(operands); $_state.addAttribute("callee", mlir::SymbolRefAttr::get(callee)); - $_state.addTypes(callee.getType().getResults()); + $_state.addTypes(callee.getFunctionType().getResults()); }]>, OpBuilder<(ins "FlatSymbolRefAttr":$callee, "ArrayRef<Type>":$resultTypes, CArg<"ValueRange", "{}">:$operands),
diff --git a/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp b/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp index 0566e5d..08d03a5 100644 --- a/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp +++ b/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp
@@ -481,7 +481,7 @@ // Generate the signature calling convention string based on types. auto cconv = makeImportCallingConventionString(importOp); if (!cconv.hasValue()) return {}; - return createFunctionSignatureDef(importOp.getType(), typeTable, + return createFunctionSignatureDef(importOp.getFunctionType(), typeTable, cconv.getValue(), /*reflectionAttrsRef=*/0, fbb); } @@ -514,7 +514,7 @@ fbb, reflectionAttrRefs.data(), reflectionAttrRefs.size()); } - return createFunctionSignatureDef(funcOp.getType(), typeTable, + return createFunctionSignatureDef(funcOp.getFunctionType(), typeTable, cconv.getValue(), reflectionAttrsRef, fbb); } @@ -525,7 +525,7 @@ // Generate the signature calling convention string based on types. auto cconv = makeCallingConventionString(funcOp); if (!cconv.hasValue()) return {}; - return createFunctionSignatureDef(funcOp.getType(), typeTable, + return createFunctionSignatureDef(funcOp.getFunctionType(), typeTable, cconv.getValue(), /*reflectionAttrsRef=*/0, fbb); }
diff --git a/iree/compiler/Dialect/VM/Target/C/CModuleTarget.cpp b/iree/compiler/Dialect/VM/Target/C/CModuleTarget.cpp index e94294b..30c0026 100644 --- a/iree/compiler/Dialect/VM/Target/C/CModuleTarget.cpp +++ b/iree/compiler/Dialect/VM/Target/C/CModuleTarget.cpp
@@ -39,12 +39,13 @@ } static LogicalResult printFunctionDeclaration( - mlir::FuncOp funcOp, llvm::raw_ostream &output, + mlir::func::FuncOp funcOp, llvm::raw_ostream &output, mlir::emitc::CppEmitter &emitter) { Operation *op = funcOp.getOperation(); if (op->hasAttr("emitc.static")) output << "static "; - if (failed(emitter.emitTypes(funcOp.getLoc(), funcOp.getType().getResults()))) + if (failed(emitter.emitTypes(funcOp.getLoc(), + funcOp.getFunctionType().getResults()))) return failure(); output << " " << funcOp.getName(); @@ -165,7 +166,7 @@ }); for (auto exportOp : exportOps) { StringRef funcName = exportOp.function_ref(); - auto funcOp = symbolTable.lookup<mlir::FuncOp>(funcName); + auto funcOp = symbolTable.lookup<mlir::func::FuncOp>(funcName); if (!funcOp) { return exportOp.emitError("Couldn't find referenced FuncOp"); } @@ -220,7 +221,7 @@ // implementation. for (auto exportOp : exportOps) { StringRef funcName = exportOp.function_ref(); - auto funcOp = symbolTable.lookup<mlir::FuncOp>(funcName); + auto funcOp = symbolTable.lookup<mlir::func::FuncOp>(funcName); if (!funcOp) { return exportOp.emitError("Couldn't find referenced FuncOp"); } @@ -362,7 +363,7 @@ output << "\n"; mlir::emitc::CppEmitter emitter(output, /*declareVariablesAtTop=*/true); - for (auto funcOp : moduleOp.getOps<mlir::FuncOp>()) { + for (auto funcOp : moduleOp.getOps<mlir::func::FuncOp>()) { Operation *op = funcOp.getOperation(); if (!op->hasAttr("vm.module.constructor")) continue; if (failed(printFunctionDeclaration(funcOp, output, emitter))) @@ -403,7 +404,7 @@ // translate functions output << "// DECLARE FUNCTIONS\n"; - for (auto funcOp : moduleOp.getOps<mlir::FuncOp>()) { + for (auto funcOp : moduleOp.getOps<mlir::func::FuncOp>()) { Operation *op = funcOp.getOperation(); if (op->hasAttr("vm.module.constructor")) continue; if (failed(printFunctionDeclaration(funcOp, output, emitter))) @@ -417,7 +418,7 @@ // TODO(simon-camp): Clean up. We generate calls to a macro that defines a // struct. As we declare all variables at the start of the function, the // macro call cannot be inlined into the function. - if (!isa<mlir::FuncOp, emitc::CallOp>(op)) continue; + if (!isa<mlir::func::FuncOp, emitc::CallOp>(op)) continue; if (op.hasAttr("vm.emit_at_end")) continue; if (op.hasAttr("emitc.static")) output << "static "; if (failed(emitter.emitOperation(op, @@ -433,7 +434,7 @@ } // Emit code for functions marked with `vm.emit_at_end`. - for (auto funcOp : moduleOp.getOps<mlir::FuncOp>()) { + for (auto funcOp : moduleOp.getOps<mlir::func::FuncOp>()) { Operation *op = funcOp.getOperation(); if (!op->hasAttr("vm.emit_at_end")) continue; if (op->hasAttr("emitc.static")) output << "static ";
diff --git a/iree/compiler/Dialect/VM/Target/C/TranslateToCpp.cpp b/iree/compiler/Dialect/VM/Target/C/TranslateToCpp.cpp index 4814217..b24c8ef 100644 --- a/iree/compiler/Dialect/VM/Target/C/TranslateToCpp.cpp +++ b/iree/compiler/Dialect/VM/Target/C/TranslateToCpp.cpp
@@ -423,7 +423,7 @@ CppEmitter::Scope scope(emitter); raw_indented_ostream &os = emitter.ostream(); if (failed(emitter.emitTypes(functionOp.getLoc(), - functionOp.getType().getResults()))) + functionOp.getFunctionType().getResults()))) return failure(); os << " " << functionOp.getName();
diff --git a/iree/compiler/Dialect/VM/Transforms/Passes.cpp b/iree/compiler/Dialect/VM/Transforms/Passes.cpp index 0daeba2..fcb3ae1 100644 --- a/iree/compiler/Dialect/VM/Transforms/Passes.cpp +++ b/iree/compiler/Dialect/VM/Transforms/Passes.cpp
@@ -12,6 +12,7 @@ #include "iree/compiler/Dialect/VM/IR/VMOps.h" #include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h" #include "mlir/Dialect/Affine/Passes.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/PassManager.h" #include "mlir/Pass/PassRegistry.h" #include "mlir/Transforms/Passes.h" @@ -23,13 +24,14 @@ void buildVMTransformPassPipeline(OpPassManager &passManager, TargetOptions targetOptions) { - passManager.addNestedPass<mlir::FuncOp>(createLoopCoalescingPass()); + passManager.addNestedPass<mlir::func::FuncOp>(createLoopCoalescingPass()); passManager.addNestedPass<IREE::Util::InitializerOp>( createLoopInvariantCodeMotionPass()); - passManager.addNestedPass<mlir::FuncOp>(createLoopInvariantCodeMotionPass()); + passManager.addNestedPass<mlir::func::FuncOp>( + createLoopInvariantCodeMotionPass()); passManager.addNestedPass<IREE::Util::InitializerOp>( createConvertSCFToCFPass()); - passManager.addNestedPass<mlir::FuncOp>(createConvertSCFToCFPass()); + passManager.addNestedPass<mlir::func::FuncOp>(createConvertSCFToCFPass()); passManager.addPass(createCanonicalizerPass()); passManager.addPass(createCSEPass());
diff --git a/iree/compiler/Dialect/VM/Utils/CallingConvention.cpp b/iree/compiler/Dialect/VM/Utils/CallingConvention.cpp index c375056..4ecbaa9 100644 --- a/iree/compiler/Dialect/VM/Utils/CallingConvention.cpp +++ b/iree/compiler/Dialect/VM/Utils/CallingConvention.cpp
@@ -76,7 +76,7 @@ Optional<std::string> makeImportCallingConventionString( IREE::VM::ImportOp importOp) { - auto functionType = importOp.getType(); + auto functionType = importOp.getFunctionType(); if (functionType.getNumInputs() == 0 && functionType.getNumResults() == 0) { return std::string("0v_v"); // Valid but empty. } @@ -114,7 +114,7 @@ } Optional<std::string> makeCallingConventionString(IREE::VM::FuncOp funcOp) { - auto functionType = funcOp.getType(); + auto functionType = funcOp.getFunctionType(); if (functionType.getNumInputs() == 0 && functionType.getNumResults() == 0) { return std::string("0v_v"); // Valid but empty. }
diff --git a/iree/compiler/InputConversion/Common/IREEImportPublic.cpp b/iree/compiler/InputConversion/Common/IREEImportPublic.cpp index aac6181..dabf718 100644 --- a/iree/compiler/InputConversion/Common/IREEImportPublic.cpp +++ b/iree/compiler/InputConversion/Common/IREEImportPublic.cpp
@@ -144,7 +144,7 @@ LogicalResult matchAndRewrite( FuncOp srcOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - FunctionType srcFuncType = srcOp.getType(); + FunctionType srcFuncType = srcOp.getFunctionType(); TypeConverter::SignatureConversion signatureConversion( srcOp.getNumArguments()); @@ -299,13 +299,13 @@ }; target.addDynamicallyLegalOp<FuncOp>([&](FuncOp funcOp) { - for (Type type : funcOp.getType().getInputs()) { + for (Type type : funcOp.getFunctionType().getInputs()) { if (isIllegalType(type)) return false; } - for (Type type : funcOp.getType().getResults()) { + for (Type type : funcOp.getFunctionType().getResults()) { if (isIllegalType(type)) return false; } - for (Block &block : funcOp.body()) { + for (Block &block : funcOp.getBody()) { for (Type type : block.getArgumentTypes()) { if (isIllegalType(type)) return false; }
diff --git a/iree/compiler/InputConversion/Common/PassDetail.h b/iree/compiler/InputConversion/Common/PassDetail.h index 71990b3..67316a7 100644 --- a/iree/compiler/InputConversion/Common/PassDetail.h +++ b/iree/compiler/InputConversion/Common/PassDetail.h
@@ -7,6 +7,7 @@ #ifndef IREE_COMPILER_INPUTCONVERSION_COMMON_PASSDETAIL_H_ #define IREE_COMPILER_INPUTCONVERSION_COMMON_PASSDETAIL_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir {
diff --git a/iree/compiler/InputConversion/Common/Passes.h b/iree/compiler/InputConversion/Common/Passes.h index a78af31..5781a19 100644 --- a/iree/compiler/InputConversion/Common/Passes.h +++ b/iree/compiler/InputConversion/Common/Passes.h
@@ -7,6 +7,7 @@ #ifndef IREE_COMPILER_INPUTCONVERSION_COMMON_PASSES_H_ #define IREE_COMPILER_INPUTCONVERSION_COMMON_PASSES_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" @@ -25,9 +26,9 @@ // Passes //===----------------------------------------------------------------------===// -std::unique_ptr<OperationPass<FuncOp>> createTopLevelSCFToCFGPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createTopLevelSCFToCFGPass(); std::unique_ptr<OperationPass<ModuleOp>> createIREEImportPublicPass(); -std::unique_ptr<OperationPass<FuncOp>> +std::unique_ptr<OperationPass<func::FuncOp>> createLinalgQuantizedMatmulToMatmulPass(); //===----------------------------------------------------------------------===//
diff --git a/iree/compiler/InputConversion/Common/test/iree_import_public.mlir b/iree/compiler/InputConversion/Common/test/iree_import_public.mlir index bc77986..b6d0670 100644 --- a/iree/compiler/InputConversion/Common/test/iree_import_public.mlir +++ b/iree/compiler/InputConversion/Common/test/iree_import_public.mlir
@@ -3,14 +3,14 @@ // CHECK-LABEL: func @bv_func // CHECK-SAME: (%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> (!hal.buffer_view, !hal.buffer_view) // CHECK: return %arg0, %arg1 : !hal.buffer_view, !hal.buffer_view -builtin.func @bv_func(%arg0 : !iree_input.buffer_view, %arg1 : !iree_input.buffer_view) -> (!iree_input.buffer_view, !iree_input.buffer_view) { +func.func @bv_func(%arg0 : !iree_input.buffer_view, %arg1 : !iree_input.buffer_view) -> (!iree_input.buffer_view, !iree_input.buffer_view) { return %arg0, %arg1 : !iree_input.buffer_view, !iree_input.buffer_view } // ----- // CHECK-LABEL: func @list_func // CHECK-SAME: (%arg0: !util.list<?>) -> !util.list<?> -builtin.func @list_func(%arg0 : !iree_input.list<!iree_input.variant>) -> !iree_input.list<!iree_input.variant> { +func.func @list_func(%arg0 : !iree_input.list<!iree_input.variant>) -> !iree_input.list<!iree_input.variant> { return %arg0 : !iree_input.list<!iree_input.variant> } @@ -18,7 +18,7 @@ // CHECK-LABEL: func @list_func_retains_iree_abi // CHECK-SAME: (%arg0: !util.list<?>) -> !util.list<?> // CHECK-SAME: iree.abi = "FOOBAR" -builtin.func @list_func_retains_iree_abi(%arg0 : !iree_input.list<!iree_input.variant>) -> !iree_input.list<!iree_input.variant> +func.func @list_func_retains_iree_abi(%arg0 : !iree_input.list<!iree_input.variant>) -> !iree_input.list<!iree_input.variant> attributes {iree.abi = "FOOBAR"} { return %arg0 : !iree_input.list<!iree_input.variant> } @@ -26,7 +26,7 @@ // ----- // CHECK-LABEL: func @list_func_call // CHECK: call @list_func_call(%arg0) : (!util.list<?>) -> !util.list<?> -builtin.func @list_func_call(%arg0 : !iree_input.list<!iree_input.variant>) -> !iree_input.list<!iree_input.variant> { +func.func @list_func_call(%arg0 : !iree_input.list<!iree_input.variant>) -> !iree_input.list<!iree_input.variant> { call @list_func_call(%arg0) : (!iree_input.list<!iree_input.variant>) -> !iree_input.list<!iree_input.variant> return %arg0 : !iree_input.list<!iree_input.variant> } @@ -34,14 +34,14 @@ // ----- // CHECK-LABEL: func @ptr_func // CHECK-SAME: (%arg0: !util.ptr<!hal.buffer_view>) -> !util.ptr<!hal.buffer_view> -builtin.func @ptr_func(%arg0 : !iree_input.ptr<!iree_input.buffer_view>) -> !iree_input.ptr<!iree_input.buffer_view> { +func.func @ptr_func(%arg0 : !iree_input.ptr<!iree_input.buffer_view>) -> !iree_input.ptr<!iree_input.buffer_view> { return %arg0 : !iree_input.ptr<!iree_input.buffer_view> } // ----- // CHECK-LABEL: func @null_op // CHECK: util.null : !util.variant -builtin.func @null_op() -> !iree_input.variant { +func.func @null_op() -> !iree_input.variant { %0 = iree_input.null : !iree_input.variant return %0 : !iree_input.variant } @@ -49,7 +49,7 @@ // ----- // CHECK-LABEL: func @tensor_to_buffer_view // CHECK: hal.tensor.export %arg0 : tensor<?x?x3xf32>{%arg1, %arg2} -> !hal.buffer_view -builtin.func @tensor_to_buffer_view(%arg0 : tensor<?x?x3xf32>, %arg1 : index, %arg2 : index) -> !iree_input.buffer_view { +func.func @tensor_to_buffer_view(%arg0 : tensor<?x?x3xf32>, %arg1 : index, %arg2 : index) -> !iree_input.buffer_view { %0 = iree_input.cast.tensor_to_buffer_view %arg0 : tensor<?x?x3xf32>{%arg1, %arg2} -> !iree_input.buffer_view return %0 : !iree_input.buffer_view } @@ -57,7 +57,7 @@ // ----- // CHECK-LABEL: func @tensor_to_buffer_view_static // CHECK: hal.tensor.export %arg0 : tensor<3xf32> -> !hal.buffer_view -builtin.func @tensor_to_buffer_view_static(%arg0 : tensor<3xf32>) -> !iree_input.buffer_view { +func.func @tensor_to_buffer_view_static(%arg0 : tensor<3xf32>) -> !iree_input.buffer_view { %0 = iree_input.cast.tensor_to_buffer_view %arg0 : tensor<3xf32> -> !iree_input.buffer_view return %0 : !iree_input.buffer_view } @@ -69,7 +69,7 @@ // CHECK: %[[ONE:.*]] = arith.constant 1 // CHECK: %[[D1:.*]] = tensor.dim %arg0, %[[ONE]] // CHECK: hal.tensor.export %arg0 : tensor<?x?x3xf32>{%[[D0]], %[[D1]]} -> !hal.buffer_view -builtin.func @tensor_to_buffer_view_implicit_dims(%arg0 : tensor<?x?x3xf32>) -> !iree_input.buffer_view { +func.func @tensor_to_buffer_view_implicit_dims(%arg0 : tensor<?x?x3xf32>) -> !iree_input.buffer_view { %0 = iree_input.cast.tensor_to_buffer_view %arg0 : tensor<?x?x3xf32> -> !iree_input.buffer_view return %0 : !iree_input.buffer_view } @@ -77,7 +77,7 @@ // ----- // CHECK-LABEL: func @buffer_view_to_tensor // CHECK: hal.tensor.import %arg0 : !hal.buffer_view -> tensor<?x?x3xf32>{%arg1, %arg2} -builtin.func @buffer_view_to_tensor(%arg0 : !iree_input.buffer_view, %arg1 : index, %arg2 : index) -> tensor<?x?x3xf32> { +func.func @buffer_view_to_tensor(%arg0 : !iree_input.buffer_view, %arg1 : index, %arg2 : index) -> tensor<?x?x3xf32> { %0 = iree_input.cast.buffer_view_to_tensor %arg0 : !iree_input.buffer_view -> tensor<?x?x3xf32>{%arg1, %arg2} return %0 : tensor<?x?x3xf32> } @@ -85,7 +85,7 @@ // ----- // CHECK-LABEL: func @buffer_view_to_tensor_static // CHECK: hal.tensor.import %arg0 : !hal.buffer_view -> tensor<3xf32> -builtin.func @buffer_view_to_tensor_static(%arg0 : !iree_input.buffer_view) -> tensor<3xf32> { +func.func @buffer_view_to_tensor_static(%arg0 : !iree_input.buffer_view) -> tensor<3xf32> { %0 = iree_input.cast.buffer_view_to_tensor %arg0 : !iree_input.buffer_view -> tensor<3xf32> return %0 : tensor<3xf32> } @@ -95,7 +95,7 @@ // CHECK: %[[D0:.*]] = hal.buffer_view.dim<%arg0 : !hal.buffer_view>[0] : index // CHECK: %[[D1:.*]] = hal.buffer_view.dim<%arg0 : !hal.buffer_view>[1] : index // CHECK: hal.tensor.import %arg0 : !hal.buffer_view -> tensor<?x?x3xf32>{%[[D0]], %[[D1]]} -builtin.func @buffer_view_to_tensor_implicit_dims(%arg0 : !iree_input.buffer_view) -> tensor<?x?x3xf32> { +func.func @buffer_view_to_tensor_implicit_dims(%arg0 : !iree_input.buffer_view) -> tensor<?x?x3xf32> { %0 = iree_input.cast.buffer_view_to_tensor %arg0 : !iree_input.buffer_view -> tensor<?x?x3xf32> return %0 : tensor<?x?x3xf32> } @@ -103,7 +103,7 @@ // ----- // CHECK-LABEL: func @buffer_view_rank // CHECK: hal.buffer_view.rank<%arg0 : !hal.buffer_view> : index -builtin.func @buffer_view_rank(%arg0 : !iree_input.buffer_view) -> index { +func.func @buffer_view_rank(%arg0 : !iree_input.buffer_view) -> index { %0 = iree_input.buffer_view.rank %arg0 : index return %0 : index } @@ -111,7 +111,7 @@ // ----- // CHECK-LABEL: func @buffer_view_dim // CHECK: hal.buffer_view.dim<%arg0 : !hal.buffer_view>[0] : index -builtin.func @buffer_view_dim(%arg0 : !iree_input.buffer_view) -> index { +func.func @buffer_view_dim(%arg0 : !iree_input.buffer_view) -> index { %0 = iree_input.buffer_view.dim %arg0, 0 : index return %0: index } @@ -119,7 +119,7 @@ // ----- // CHECK-LABEL: func @list_create // CHECK: util.list.create %arg0 : !util.list<?> -builtin.func @list_create(%arg0 : index) -> !iree_input.list<!iree_input.variant> { +func.func @list_create(%arg0 : index) -> !iree_input.list<!iree_input.variant> { %0 = iree_input.list.create %arg0 : !iree_input.list<!iree_input.variant> return %0 : !iree_input.list<!iree_input.variant> } @@ -127,7 +127,7 @@ // ----- // CHECK-LABEL: func @list_size // CHECK: util.list.size %arg0 : !util.list<?> -builtin.func @list_size(%arg0 : !iree_input.list<!iree_input.variant>) -> index { +func.func @list_size(%arg0 : !iree_input.list<!iree_input.variant>) -> index { %0 = iree_input.list.size %arg0 : !iree_input.list<!iree_input.variant> return %0 : index } @@ -135,7 +135,7 @@ // ----- // CHECK-LABEL: func @list_resize // CHECK: util.list.resize %arg0, %arg1 : !util.list<?> -builtin.func @list_resize(%arg0 : !iree_input.list<!iree_input.variant>, %arg1 : index) { +func.func @list_resize(%arg0 : !iree_input.list<!iree_input.variant>, %arg1 : index) { iree_input.list.resize %arg0, %arg1 : !iree_input.list<!iree_input.variant> return } @@ -143,7 +143,7 @@ // ----- // CHECK-LABEL: func @list_get // CHECK: util.list.get %arg0[%arg1] : !util.list<?> -builtin.func @list_get(%arg0 : !iree_input.list<!iree_input.variant>, %arg1 : index) -> !iree_input.variant { +func.func @list_get(%arg0 : !iree_input.list<!iree_input.variant>, %arg1 : index) -> !iree_input.variant { %0 = iree_input.list.get %arg0[%arg1] : !iree_input.list<!iree_input.variant> -> !iree_input.variant return %0 : !iree_input.variant } @@ -151,7 +151,7 @@ // ----- // CHECK-LABEL: func @list_set // CHECK: util.list.set %arg0[%arg1], %arg2 : !util.list<?> -builtin.func @list_set(%arg0 : !iree_input.list<!iree_input.variant>, %arg1 : index, %arg2 : !iree_input.variant) { +func.func @list_set(%arg0 : !iree_input.list<!iree_input.variant>, %arg1 : index, %arg2 : !iree_input.variant) { iree_input.list.set %arg0[%arg1], %arg2 : !iree_input.list<!iree_input.variant>, !iree_input.variant return } @@ -159,7 +159,7 @@ // ----- // CHECK-LABEL: func @tensor_reshape // CHECK: flow.tensor.reshape %arg0 : tensor<?x?xf32>{%arg1, %arg2} -> tensor<?x?xf32>{%arg2, %arg1} -builtin.func @tensor_reshape(%arg0 : tensor<?x?xf32>, %arg1 : index, %arg2 : index) -> tensor<?x?xf32> { +func.func @tensor_reshape(%arg0 : tensor<?x?xf32>, %arg1 : index, %arg2 : index) -> tensor<?x?xf32> { %0 = iree_input.tensor.reshape %arg0 : tensor<?x?xf32>{%arg1, %arg2} -> tensor<?x?xf32>{%arg2, %arg1} return %0 : tensor<?x?xf32> } @@ -167,7 +167,7 @@ // ----- // CHECK-LABEL: func @tensor_load // CHECK: flow.tensor.load %arg0[%arg2, %arg3] : tensor<?x3xf32>{%arg1} -builtin.func @tensor_load(%arg0 : tensor<?x3xf32>, %arg1 : index, %arg2 : index, %arg3 : index) -> f32 { +func.func @tensor_load(%arg0 : tensor<?x3xf32>, %arg1 : index, %arg2 : index, %arg3 : index) -> f32 { %0 = iree_input.tensor.load %arg0[%arg2, %arg3] : tensor<?x3xf32>{%arg1} return %0 : f32 } @@ -175,7 +175,7 @@ // ----- // CHECK-LABEL: func @tensor_store // CHECK: flow.tensor.store %arg4, %arg0[%arg2, %arg3] : tensor<?x3xf32>{%arg1} -builtin.func @tensor_store(%arg0 : tensor<?x3xf32>, %arg1 : index, %arg2 : index, %arg3 : index, %arg4 : f32) { +func.func @tensor_store(%arg0 : tensor<?x3xf32>, %arg1 : index, %arg2 : index, %arg3 : index, %arg4 : f32) { iree_input.tensor.store %arg4, %arg0[%arg2, %arg3] : tensor<?x3xf32>{%arg1} return } @@ -183,7 +183,7 @@ // ----- // CHECK-LABEL: func @tensor_splat // CHECK: flow.tensor.splat %arg0 : tensor<?x?xf32>{%arg1, %arg2} -builtin.func @tensor_splat(%arg0 : f32, %arg1 : index, %arg2 : index) -> tensor<?x?xf32> { +func.func @tensor_splat(%arg0 : f32, %arg1 : index, %arg2 : index) -> tensor<?x?xf32> { %0 = iree_input.tensor.splat %arg0 : tensor<?x?xf32>{%arg1, %arg2} return %0 : tensor<?x?xf32> } @@ -191,7 +191,7 @@ // ----- // CHECK-LABEL: func @tensor_clone // CHECK: flow.tensor.clone %arg0 : tensor<?x?xf32>{%arg1, %arg2} -builtin.func @tensor_clone(%arg0 : tensor<?x?xf32>, %arg1 : index, %arg2 : index) -> tensor<?x?xf32> { +func.func @tensor_clone(%arg0 : tensor<?x?xf32>, %arg1 : index, %arg2 : index) -> tensor<?x?xf32> { %0 = iree_input.tensor.clone %arg0 : tensor<?x?xf32>{%arg1, %arg2} return %0 : tensor<?x?xf32> } @@ -199,7 +199,7 @@ // ----- // CHECK-LABEL: func @tensor_slice // CHECK: flow.tensor.slice %arg0[%arg1 for %arg2] : tensor<?xf32>{%arg3} -> tensor<?xf32>{%arg4} -builtin.func @tensor_slice(%arg0 : tensor<?xf32>, %arg1 : index, %arg2 : index, %arg3 : index, %arg4 : index) -> tensor<?xf32> { +func.func @tensor_slice(%arg0 : tensor<?xf32>, %arg1 : index, %arg2 : index, %arg3 : index, %arg4 : index) -> tensor<?xf32> { %0 = iree_input.tensor.slice %arg0[%arg1 for %arg2] : tensor<?xf32>{%arg3} -> tensor<?xf32>{%arg4} return %0 : tensor<?xf32> } @@ -207,7 +207,7 @@ // ----- // CHECK-LABEL: func @tensor_update // CHECK: flow.tensor.update %arg3, %arg0[%arg1] : tensor<?xf32>{%arg2} -> %arg0 as tensor<?xf32>{%arg4} -builtin.func @tensor_update(%arg0 : tensor<?xf32>, %arg1 : index, %arg2 : index, %arg3 : tensor<?xf32>, %arg4 : index) -> tensor<?xf32> { +func.func @tensor_update(%arg0 : tensor<?xf32>, %arg1 : index, %arg2 : index, %arg3 : tensor<?xf32>, %arg4 : index) -> tensor<?xf32> { %0 = iree_input.tensor.update %arg3, %arg0[%arg1] : tensor<?xf32>{%arg2} -> tensor<?xf32>{%arg4} return %0 : tensor<?xf32> } @@ -215,7 +215,7 @@ // ----- // CHECK-LABEL: func @tensor_trace // CHECK: flow.tensor.trace {key = "FOOBAR"} %arg0, %arg1 : tensor<5xf32>, tensor<3xf32> -builtin.func @tensor_trace(%arg0 : tensor<5xf32>, %arg1 : tensor<3xf32>) { +func.func @tensor_trace(%arg0 : tensor<5xf32>, %arg1 : tensor<3xf32>) { iree_input.tensor.trace "FOOBAR" %arg0, %arg1 : tensor<5xf32>, tensor<3xf32> return } @@ -240,7 +240,7 @@ // CHECK-NEXT: util.initializer.return // CHECK-NEXT: } // CHECK: func private @initializer() -> tensor<4xi32> - builtin.func private @initializer() -> tensor<4xi32> + func.func private @initializer() -> tensor<4xi32> } // -----
diff --git a/iree/compiler/InputConversion/MHLO/ConvertComplexToReal.cpp b/iree/compiler/InputConversion/MHLO/ConvertComplexToReal.cpp index 09c2ab4..35bbb18 100644 --- a/iree/compiler/InputConversion/MHLO/ConvertComplexToReal.cpp +++ b/iree/compiler/InputConversion/MHLO/ConvertComplexToReal.cpp
@@ -227,7 +227,7 @@ ConvertCompareOp(TypeConverter &typeConverter, MLIRContext *context, mhlo::ComparisonDirection direction) : OpConversionPattern<CompareOpTy>(typeConverter, context), - direction(mhlo::stringifyEnum(direction)) {} + direction(direction) {} LogicalResult matchAndRewrite( CompareOpTy op, typename CompareOpTy::Adaptor adaptor, @@ -262,7 +262,7 @@ return success(); } - StringRef direction; + mhlo::ComparisonDirection direction; }; struct ElideComplexPattern : public OpConversionPattern<mhlo::ComplexOp> {
diff --git a/iree/compiler/InputConversion/MHLO/FlattenTuplesInCFG.cpp b/iree/compiler/InputConversion/MHLO/FlattenTuplesInCFG.cpp index 23eabe9..5d2bddf 100644 --- a/iree/compiler/InputConversion/MHLO/FlattenTuplesInCFG.cpp +++ b/iree/compiler/InputConversion/MHLO/FlattenTuplesInCFG.cpp
@@ -288,7 +288,7 @@ std::vector<std::pair<FuncOp, FuncOp>> convertedFunctions; for (auto oldFunction : module.getOps<FuncOp>()) { - auto oldFunctionType = oldFunction.getType(); + auto oldFunctionType = oldFunction.getFunctionType(); llvm::SmallVector<Type, 10> newInputTypes; untupleTypes(oldFunctionType.getInputs(), &newInputTypes);
diff --git a/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp b/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp index 04264ae..ac40eca 100644 --- a/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp +++ b/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp
@@ -181,10 +181,10 @@ return success(); } -static LogicalResult convertFunc(mlir::FuncOp oldFuncOp, +static LogicalResult convertFunc(mlir::func::FuncOp oldFuncOp, FlowTypeConverter &typeConverter, OpBuilder &moduleBuilder) { - auto oldType = oldFuncOp.getType(); + auto oldType = oldFuncOp.getFunctionType(); TypeConverter::SignatureConversion signature(oldType.getNumInputs()); for (unsigned i = 0, e = oldType.getNumInputs(); i != e; ++i) { if (failed(typeConverter.convertSignatureArg(i, oldType.getInput(i), @@ -224,7 +224,7 @@ for (auto *oldOp : oldOps) { OpBuilder moduleBuilder(moduleOp); moduleBuilder.setInsertionPoint(oldOp); - if (auto oldFuncOp = dyn_cast<mlir::FuncOp>(oldOp)) { + if (auto oldFuncOp = dyn_cast<mlir::func::FuncOp>(oldOp)) { if (failed(convertFunc(oldFuncOp, typeConverter, moduleBuilder))) { return signalPassFailure(); }
diff --git a/iree/compiler/InputConversion/MHLO/MHLOToLinalgOnTensors.cpp b/iree/compiler/InputConversion/MHLO/MHLOToLinalgOnTensors.cpp index ebc4504..3e94508 100644 --- a/iree/compiler/InputConversion/MHLO/MHLOToLinalgOnTensors.cpp +++ b/iree/compiler/InputConversion/MHLO/MHLOToLinalgOnTensors.cpp
@@ -178,7 +178,7 @@ LogicalResult matchAndRewrite( mhlo::FftOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - if (op.fft_type() != "RFFT") { + if (op.fft_type() != mhlo::FftType::RFFT) { return rewriter.notifyMatchFailure(op, "non RFFT types are supported yet"); } @@ -222,7 +222,7 @@ LogicalResult matchAndRewrite( FuncOp srcOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - FunctionType srcFuncType = srcOp.getType(); + FunctionType srcFuncType = srcOp.getFunctionType(); TypeConverter::SignatureConversion signatureConversion( srcOp.getNumArguments()); @@ -349,13 +349,13 @@ // Functions must have legal types. target.addDynamicallyLegalOp<FuncOp>([&](FuncOp funcOp) { - for (Type type : funcOp.getType().getInputs()) { + for (Type type : funcOp.getFunctionType().getInputs()) { if (isIllegalType(type)) return false; } - for (Type type : funcOp.getType().getResults()) { + for (Type type : funcOp.getFunctionType().getResults()) { if (isIllegalType(type)) return false; } - for (Block &block : funcOp.body()) { + for (Block &block : funcOp.getBody()) { for (Type type : block.getArgumentTypes()) { if (isIllegalType(type)) return false; }
diff --git a/iree/compiler/InputConversion/MHLO/PassDetail.h b/iree/compiler/InputConversion/MHLO/PassDetail.h index f7098ac..097b977 100644 --- a/iree/compiler/InputConversion/MHLO/PassDetail.h +++ b/iree/compiler/InputConversion/MHLO/PassDetail.h
@@ -7,6 +7,7 @@ #ifndef IREE_COMPILER_INPUTCONVERSION_MHLO_PASSDETAIL_H_ #define IREE_COMPILER_INPUTCONVERSION_MHLO_PASSDETAIL_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir {
diff --git a/iree/compiler/InputConversion/MHLO/Passes.h b/iree/compiler/InputConversion/MHLO/Passes.h index a027fc3..9c82ece 100644 --- a/iree/compiler/InputConversion/MHLO/Passes.h +++ b/iree/compiler/InputConversion/MHLO/Passes.h
@@ -7,6 +7,7 @@ #ifndef IREE_COMPILER_INPUTCONVERSION_MHLO_PASSES_H_ #define IREE_COMPILER_INPUTCONVERSION_MHLO_PASSES_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir { @@ -52,15 +53,16 @@ std::unique_ptr<OperationPass<ModuleOp>> createLegalizeInputTypesPass(); /// Creates XLA-HLO to Linalg on tensors transformation pass. -std::unique_ptr<OperationPass<FuncOp>> createMHLOToLinalgOnTensorsPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createMHLOToLinalgOnTensorsPass(); /// Creates XLA-HLO to LinalgExt pass. -std::unique_ptr<OperationPass<FuncOp>> createConvertMHLOToLinalgExtPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createConvertMHLOToLinalgExtPass(); /// Creates XLA-HLO preprocessing transformation pass. In this pass we should /// have all mhlo -> mhlo transformations that are shared between all /// backends. -std::unique_ptr<OperationPass<FuncOp>> createMHLOToMHLOPreprocessingPass(); +std::unique_ptr<OperationPass<func::FuncOp>> +createMHLOToMHLOPreprocessingPass(); // Verifies a module being input to the core compiler pipeline only contains // IR structures that are supported at that level. @@ -71,7 +73,8 @@ // Test passes //------------------------------------------------------------------------------ -std::unique_ptr<OperationPass<FuncOp>> createTestMHLOConvertComplexToRealPass(); +std::unique_ptr<OperationPass<func::FuncOp>> +createTestMHLOConvertComplexToRealPass(); //===----------------------------------------------------------------------===// // Register all Passes
diff --git a/iree/compiler/InputConversion/MHLO/Passes.td b/iree/compiler/InputConversion/MHLO/Passes.td index 4cf16ae..3960c0d 100644 --- a/iree/compiler/InputConversion/MHLO/Passes.td +++ b/iree/compiler/InputConversion/MHLO/Passes.td
@@ -10,13 +10,13 @@ include "mlir/Pass/PassBase.td" def ConvertMHLOToLinalgOnTensors : - Pass<"iree-mhlo-to-linalg-on-tensors", "FuncOp"> { + Pass<"iree-mhlo-to-linalg-on-tensors", "func::FuncOp"> { let summary = "Convert from XLA-HLO ops to Linalg ops on tensors"; let constructor = "mlir::iree_compiler::MHLO::createMHLOToLinalgOnTensorsPass()"; } def ConvertMHLOToLinalgExt - : Pass<"iree-mhlo-to-linalg-ext", "FuncOp"> { + : Pass<"iree-mhlo-to-linalg-ext", "func::FuncOp"> { let summary = "Convert from XLA-HLO ops to LinalgExt ops and distribute to Flow ops"; let constructor = @@ -36,7 +36,7 @@ } def MHLOToMHLOPreprocessing : - Pass<"iree-mhlo-to-mhlo-preprocessing", "FuncOp"> { + Pass<"iree-mhlo-to-mhlo-preprocessing", "func::FuncOp"> { let summary = "Apply mhlo to mhlo transformations for some mhlo ops"; let constructor = "mlir::iree_compiler::MHLO::createMHLOToMHLOPreprocessingPass()"; let options = [ @@ -56,7 +56,7 @@ //------------------------------------------------------------------------------ def TestMHLOConvertComplexToReal : - Pass<"iree-test-mhlo-convert-complex-to-real", "FuncOp"> { + Pass<"iree-test-mhlo-convert-complex-to-real", "func::FuncOp"> { let summary = "Test pass that does an MHLO->MHLO conversion of just complex arithmetic ops."; let constructor = "mlir::iree_compiler::MHLO::createTestMHLOConvertComplexToRealPass()"; }
diff --git a/iree/compiler/InputConversion/MHLO/test/broadcasting.mlir b/iree/compiler/InputConversion/MHLO/test/broadcasting.mlir index 8999445..e3b8bc9 100644 --- a/iree/compiler/InputConversion/MHLO/test/broadcasting.mlir +++ b/iree/compiler/InputConversion/MHLO/test/broadcasting.mlir
@@ -79,7 +79,7 @@ // NOTE: compare is unique because of the element type switch. The pattern // will fail or the verifier will catch it if wrong. // CHECK-NOT: mhlo.compare - %0 = chlo.broadcast_compare %arg0, %arg1 {comparison_direction = "EQ"} : (tensor<?xf32>, tensor<?x?xf32>) -> tensor<?x?xi1> + %0 = chlo.broadcast_compare %arg0, %arg1 {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<?xf32>, tensor<?x?xf32>) -> tensor<?x?xi1> return %0 : tensor<?x?xi1> } @@ -296,7 +296,7 @@ // CHECK-LABEL: @compareWithoutBroadcast func @compareWithoutBroadcast(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xi1> { // CHECK-NOT: mhlo.compare - %0 = chlo.broadcast_compare %arg0, %arg1 {comparison_direction = "EQ"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xi1> + %0 = chlo.broadcast_compare %arg0, %arg1 {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xi1> return %0 : tensor<4xi1> }
diff --git a/iree/compiler/InputConversion/MHLO/test/convert_complex_to_real.mlir b/iree/compiler/InputConversion/MHLO/test/convert_complex_to_real.mlir index 68ae132..b3c4d4c 100644 --- a/iree/compiler/InputConversion/MHLO/test/convert_complex_to_real.mlir +++ b/iree/compiler/InputConversion/MHLO/test/convert_complex_to_real.mlir
@@ -123,10 +123,10 @@ %arg2 : tensor<2xf32>, %arg3 : tensor<2xf32>) -> (tensor<2xi1>) { %lhs = "mhlo.complex"(%arg0, %arg1) : (tensor<2xf32>, tensor<2xf32>) -> (tensor<2xcomplex<f32>>) %rhs = "mhlo.complex"(%arg2, %arg3) : (tensor<2xf32>, tensor<2xf32>) -> (tensor<2xcomplex<f32>>) - // CHECK-DAG: %[[OUTR:.+]] = chlo.broadcast_compare %arg0, %arg2 {comparison_direction = "EQ"} - // CHECK-DAG: %[[OUTI:.+]] = chlo.broadcast_compare %arg1, %arg3 {comparison_direction = "EQ"} + // CHECK-DAG: %[[OUTR:.+]] = chlo.broadcast_compare %arg0, %arg2 {comparison_direction = #mhlo<"comparison_direction EQ">} + // CHECK-DAG: %[[OUTI:.+]] = chlo.broadcast_compare %arg1, %arg3 {comparison_direction = #mhlo<"comparison_direction EQ">} // CHECK-DAG: %[[OUT:.+]] = mhlo.and %[[OUTR]], %[[OUTI]] - %0 = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<2xcomplex<f32>>, tensor<2xcomplex<f32>>) -> tensor<2xi1> + %0 = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<2xcomplex<f32>>, tensor<2xcomplex<f32>>) -> tensor<2xi1> // CHECK: return %[[OUT]] return %0 : tensor<2xi1> @@ -137,10 +137,10 @@ %arg2 : tensor<2xf32>, %arg3 : tensor<2xf32>) -> (tensor<2xi1>) { %lhs = "mhlo.complex"(%arg0, %arg1) : (tensor<2xf32>, tensor<2xf32>) -> (tensor<2xcomplex<f32>>) %rhs = "mhlo.complex"(%arg2, %arg3) : (tensor<2xf32>, tensor<2xf32>) -> (tensor<2xcomplex<f32>>) - // CHECK-DAG: %[[OUTR:.+]] = chlo.broadcast_compare %arg0, %arg2 {comparison_direction = "NE"} - // CHECK-DAG: %[[OUTI:.+]] = chlo.broadcast_compare %arg1, %arg3 {comparison_direction = "NE"} + // CHECK-DAG: %[[OUTR:.+]] = chlo.broadcast_compare %arg0, %arg2 {comparison_direction = #mhlo<"comparison_direction NE">} + // CHECK-DAG: %[[OUTI:.+]] = chlo.broadcast_compare %arg1, %arg3 {comparison_direction = #mhlo<"comparison_direction NE">} // CHECK-DAG: %[[OUT:.+]] = mhlo.or %[[OUTR]], %[[OUTI]] - %0 = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "NE"} : (tensor<2xcomplex<f32>>, tensor<2xcomplex<f32>>) -> tensor<2xi1> + %0 = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction NE">} : (tensor<2xcomplex<f32>>, tensor<2xcomplex<f32>>) -> tensor<2xi1> // CHECK: return %[[OUT]] return %0 : tensor<2xi1>
diff --git a/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir b/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir index bb2349e..8600141 100644 --- a/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir +++ b/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir
@@ -5,7 +5,7 @@ func @sort_1d(%arg0: tensor<128xi32>) -> (tensor<128xi32>) { %0 = "mhlo.sort"(%arg0) ( { ^bb0(%arg2: tensor<i32>, %arg3: tensor<i32>): - %1 = "mhlo.compare"(%arg2, %arg3) {comparison_direction = "GT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %1 = "mhlo.compare"(%arg2, %arg3) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> "mhlo.return"(%1) : (tensor<i1>) -> () }) {dimension = 0 : i64, is_stable = false} : (tensor<128xi32>) -> (tensor<128xi32>) return %0 : tensor<128xi32> @@ -26,7 +26,7 @@ func @sort_1d_ui(%arg0: tensor<128xui32>) -> (tensor<128xui32>) { %0 = "mhlo.sort"(%arg0) ( { ^bb0(%arg2: tensor<ui32>, %arg3: tensor<ui32>): // no predecessors - %1 = "mhlo.compare"(%arg2, %arg3) {comparison_direction = "GT"} : (tensor<ui32>, tensor<ui32>) -> tensor<i1> + %1 = "mhlo.compare"(%arg2, %arg3) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<ui32>, tensor<ui32>) -> tensor<i1> "mhlo.return"(%1) : (tensor<i1>) -> () }) {dimension = 0 : i64, is_stable = false} : (tensor<128xui32>) -> (tensor<128xui32>) return %0 : tensor<128xui32> @@ -50,7 +50,7 @@ %0 = mhlo.constant dense<0> : tensor<i32> %1 = "mhlo.sort"(%arg0) ( { ^bb0(%arg1: tensor<i32>, %arg3: tensor<i32>): - %2 = "mhlo.compare"(%arg1, %0) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %2 = "mhlo.compare"(%arg1, %0) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> "mhlo.return"(%2) : (tensor<i1>) -> () }) {dimension = 1 : i64, is_stable = true} : (tensor<1x10xi32>) -> tensor<1x10xi32> return %1 : tensor<1x10xi32> @@ -72,7 +72,7 @@ func @sort_argument_capture(%arg0: tensor<1x10xi32>, %arg1 : tensor<i32>) -> tensor<1x10xi32> { %1 = "mhlo.sort"(%arg0) ( { ^bb0(%arg2: tensor<i32>, %arg3: tensor<i32>): - %2 = "mhlo.compare"(%arg2, %arg1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %2 = "mhlo.compare"(%arg2, %arg1) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> "mhlo.return"(%2) : (tensor<i1>) -> () }) {dimension = 1 : i64, is_stable = true} : (tensor<1x10xi32>) -> tensor<1x10xi32> return %1 : tensor<1x10xi32> @@ -95,7 +95,7 @@ func @sort_2d(%arg0: tensor<16x32xi32>) -> (tensor<16x32xi32>) { %0 = "mhlo.sort"(%arg0) ( { ^bb0(%arg2: tensor<i32>, %arg3: tensor<i32>): - %1 = "mhlo.compare"(%arg2, %arg3) {comparison_direction = "GT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %1 = "mhlo.compare"(%arg2, %arg3) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> "mhlo.return"(%1) : (tensor<i1>) -> () }) {dimension = 0 : i64, is_stable = false} : (tensor<16x32xi32>) -> (tensor<16x32xi32>) return %0 : tensor<16x32xi32> @@ -118,7 +118,7 @@ ^bb0(%arg1: tensor<f32>, %arg2: tensor<f32>): %2 = "mhlo.bitcast_convert"(%arg1) : (tensor<f32>) -> tensor<ui32> %3 = "mhlo.bitcast_convert"(%arg2) : (tensor<f32>) -> tensor<ui32> - %4 = "mhlo.compare"(%2, %3) {comparison_direction = "LT"} : (tensor<ui32>, tensor<ui32>) -> tensor<i1> + %4 = "mhlo.compare"(%2, %3) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<ui32>, tensor<ui32>) -> tensor<i1> "mhlo.return"(%4) : (tensor<i1>) -> () }) {dimension = 1 : i64, is_stable = true} : (tensor<1x5xf32>) -> tensor<1x5xf32> return %1 : tensor<1x5xf32> @@ -144,7 +144,7 @@ %1 = "mhlo.sort"(%arg0) ( { ^bb0(%arg1: tensor<f32>, %arg2: tensor<f32>): %2 = "mhlo.bitcast_convert"(%arg1) : (tensor<f32>) -> tensor<ui32> - %3 = "mhlo.compare"(%2, %ui32) {comparison_direction = "LT"} : (tensor<ui32>, tensor<ui32>) -> tensor<i1> + %3 = "mhlo.compare"(%2, %ui32) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<ui32>, tensor<ui32>) -> tensor<i1> "mhlo.return"(%3) : (tensor<i1>) -> () }) {dimension = 1 : i64, is_stable = true} : (tensor<1x5xf32>) -> tensor<1x5xf32> return %1 : tensor<1x5xf32> @@ -176,7 +176,7 @@ %3 = mhlo.add %2, %arg1 : tensor<complex<f32>> %4 = "mhlo.real"(%3) : (tensor<complex<f32>>) -> tensor<f32> %5 = "mhlo.imag"(%3) : (tensor<complex<f32>>) -> tensor<f32> - %6 = "mhlo.compare"(%4, %5) {comparison_direction = "LT"} : (tensor<f32>, tensor<f32>) -> tensor<i1> + %6 = "mhlo.compare"(%4, %5) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<f32>, tensor<f32>) -> tensor<i1> "mhlo.return"(%6) : (tensor<i1>) -> () }) {dimension = 1 : i64, is_stable = true} : (tensor<1x5xf32>) -> tensor<1x5xf32> return %1 : tensor<1x5xf32> @@ -200,7 +200,7 @@ func @topk(%arg0: tensor<128xi32>, %arg1: tensor<128xi32>) -> (tensor<128xi32>) { %0:2 = "mhlo.sort"(%arg0, %arg1) ( { ^bb0(%arg2: tensor<i32>, %arg3: tensor<i32>, %arg4: tensor<i32>, %arg5: tensor<i32>): - %1 = "mhlo.compare"(%arg2, %arg3) {comparison_direction = "GT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %1 = "mhlo.compare"(%arg2, %arg3) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> "mhlo.return"(%1) : (tensor<i1>) -> () }) {dimension = 0 : i64, is_stable = false} : (tensor<128xi32>, tensor<128xi32>) -> (tensor<128xi32>, tensor<128xi32>) return %0#0 : tensor<128xi32> @@ -413,7 +413,7 @@ func @rfft_1d(%input: tensor<8xf32>) -> (tensor<5xf32>, tensor<5xf32>) { %0 = "mhlo.fft"(%input) { - fft_length = dense<8> : tensor<1xi64>, fft_type = "RFFT" + fft_length = dense<8> : tensor<1xi64>, fft_type = #mhlo<"fft_type RFFT"> } : (tensor<8xf32>) -> tensor<5xcomplex<f32>> %1 = "mhlo.real"(%0) : (tensor<5xcomplex<f32>>) -> tensor<5xf32> %2 = "mhlo.imag"(%0) : (tensor<5xcomplex<f32>>) -> tensor<5xf32> @@ -460,7 +460,7 @@ func @rfft_2d(%input: tensor<4x8xf32>) -> (tensor<4x5xf32>, tensor<4x5xf32>) { %0 = "mhlo.fft"(%input) { - fft_length = dense<8> : tensor<1xi64>, fft_type = "RFFT" + fft_length = dense<8> : tensor<1xi64>, fft_type = #mhlo<"fft_type RFFT"> } : (tensor<4x8xf32>) -> tensor<4x5xcomplex<f32>> %1 = "mhlo.real"(%0) : (tensor<4x5xcomplex<f32>>) -> tensor<4x5xf32> %2 = "mhlo.imag"(%0) : (tensor<4x5xcomplex<f32>>) -> tensor<4x5xf32>
diff --git a/iree/compiler/InputConversion/MHLO/test/fft.mlir b/iree/compiler/InputConversion/MHLO/test/fft.mlir index 367269b..a09bee1 100644 --- a/iree/compiler/InputConversion/MHLO/test/fft.mlir +++ b/iree/compiler/InputConversion/MHLO/test/fft.mlir
@@ -2,7 +2,7 @@ func @rfft_1d(%input: tensor<32xf32>) -> (tensor<17xf32>, tensor<17xf32>) { %0 = "mhlo.fft"(%input) { - fft_length = dense<32> : tensor<1xi64>, fft_type = "RFFT" + fft_length = dense<32> : tensor<1xi64>, fft_type = #mhlo<"fft_type RFFT"> } : (tensor<32xf32>) -> tensor<17xcomplex<f32>> %1 = "mhlo.real"(%0) : (tensor<17xcomplex<f32>>) -> tensor<17xf32> %2 = "mhlo.imag"(%0) : (tensor<17xcomplex<f32>>) -> tensor<17xf32> @@ -33,7 +33,7 @@ func @rfft_2d(%input: tensor<1x32xf32>) -> (tensor<1x17xf32>, tensor<1x17xf32>) { %0 = "mhlo.fft"(%input) { - fft_length = dense<32> : tensor<1xi64>, fft_type = "RFFT" + fft_length = dense<32> : tensor<1xi64>, fft_type = #mhlo<"fft_type RFFT"> } : (tensor<1x32xf32>) -> tensor<1x17xcomplex<f32>> %1 = "mhlo.real"(%0) : (tensor<1x17xcomplex<f32>>) -> tensor<1x17xf32> %2 = "mhlo.imag"(%0) : (tensor<1x17xcomplex<f32>>) -> tensor<1x17xf32>
diff --git a/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir b/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir index ed95270..13c3c1a 100644 --- a/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir +++ b/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir
@@ -88,7 +88,7 @@ } // ----- -// expected-error@+1 {{'builtin.func' op unable to legalize type of input 0}} +// expected-error@+1 {{'func.func' op unable to legalize type of input 0}} func @tensorUnrankedArg(%arg0 : tensor<*xi64>) -> tensor<*xi64> { return %arg0 : tensor<*xi64> } @@ -106,14 +106,14 @@ // CHECK-LABEL: func @compareI64 // CHECK-SAME: (%arg0: tensor<i32>, %arg1: tensor<i32>) -> (i1, tensor<i32>) func @compareI64(%arg0 : tensor<i64>, %arg1 : tensor<i64>) -> (i1, tensor<i64>) { - // CHECK-NEXT: %0 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + // CHECK-NEXT: %0 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> // CHECK-NEXT: %1 = tensor.extract %0[] : tensor<i1> // CHECK-NEXT: cf.cond_br %1, ^bb1(%1, %arg0 : i1, tensor<i32>), ^bb2(%1, %arg1 : i1, tensor<i32>) // CHECK-NEXT: ^bb1(%2: i1, %3: tensor<i32>): // pred: ^bb0 // CHECK-NEXT: return %2, %3 : i1, tensor<i32> // CHECK-NEXT: ^bb2(%4: i1, %5: tensor<i32>): // pred: ^bb0 // CHECK-NEXT: return %4, %5 : i1, tensor<i32> - %0 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<i64>, tensor<i64>) -> tensor<i1> + %0 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<i64>, tensor<i64>) -> tensor<i1> %1 = tensor.extract %0[] : tensor<i1> cf.cond_br %1, ^bb1(%1, %arg0 : i1, tensor<i64>), ^bb2(%1, %arg1 : i1, tensor<i64>) ^bb1(%2 : i1, %3 : tensor<i64>): @@ -161,7 +161,7 @@ // CHECK: util.global.load @[[VAR]] // CHECK: util.global.store %{{.+}}, @[[VAR]] util.global mutable @readwritevar = dense<0> : tensor<i64> -builtin.func @foo(%arg0 : tensor<i64>) { +func.func @foo(%arg0 : tensor<i64>) { %0 = util.global.load @readwritevar : tensor<i64> %1 = chlo.broadcast_add %0, %arg0 : (tensor<i64>, tensor<i64>) -> tensor<i64> util.global.store %1, @readwritevar : tensor<i64>
diff --git a/iree/compiler/InputConversion/MHLO/test/mhlo_to_mhlo_preprocessing_canoncalize_dot_general.mlir b/iree/compiler/InputConversion/MHLO/test/mhlo_to_mhlo_preprocessing_canoncalize_dot_general.mlir index fc51c87..6203cae 100644 --- a/iree/compiler/InputConversion/MHLO/test/mhlo_to_mhlo_preprocessing_canoncalize_dot_general.mlir +++ b/iree/compiler/InputConversion/MHLO/test/mhlo_to_mhlo_preprocessing_canoncalize_dot_general.mlir
@@ -8,7 +8,7 @@ rhs_batching_dimensions = [], rhs_contracting_dimensions = [0, 1], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<1x32x128x4xf32>, tensor<128x4x8x64xf32>) -> tensor<1x32x8x64xf32> return %0 : tensor<1x32x8x64xf32> } @@ -30,7 +30,7 @@ rhs_batching_dimensions = [0, 1], rhs_contracting_dimensions = [2], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<1x8x32x64xf32>, tensor<1x8x64x32xf32>) -> tensor<1x8x32x32xf32> return %0 : tensor<1x8x32x32xf32> } @@ -51,7 +51,7 @@ rhs_batching_dimensions = [0, 1], rhs_contracting_dimensions = [2], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<1x8x64x32xf32>, tensor<1x8x64x32xf32>) -> tensor<1x8x32x32xf32> return %0 : tensor<1x8x32x32xf32> } @@ -72,7 +72,7 @@ rhs_batching_dimensions = [0, 1], rhs_contracting_dimensions = [3], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<1x8x32x64xf32>, tensor<1x8x32x64xf32>) -> tensor<1x8x32x32xf32> return %0 : tensor<1x8x32x32xf32> } @@ -95,7 +95,7 @@ rhs_batching_dimensions = [0, 1], rhs_contracting_dimensions = [3], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<1x8x64x32xf32>, tensor<1x8x32x64xf32>) -> tensor<1x8x32x32xf32> return %0 : tensor<1x8x32x32xf32> } @@ -118,7 +118,7 @@ rhs_batching_dimensions = [0, 2], rhs_contracting_dimensions = [3], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<1x1x8x64xf32>, tensor<1x512x8x64xf32>) -> tensor<1x8x1x512xf32> return %0 : tensor<1x8x1x512xf32> }
diff --git a/iree/test/e2e/models/collatz.mlir b/iree/test/e2e/models/collatz.mlir index 5de052d..12b7394 100644 --- a/iree/test/e2e/models/collatz.mlir +++ b/iree/test/e2e/models/collatz.mlir
@@ -10,20 +10,20 @@ %3 = mhlo.constant dense<0.0> : tensor<f32> cf.br ^bb1(%3, %arg0 : tensor<f32>, tensor<f32>) ^bb1(%4: tensor<f32>, %5: tensor<f32>): - %6 = "mhlo.compare"(%5, %0) {comparison_direction = "GT"} : (tensor<f32>, tensor<f32>) -> tensor<i1> + %6 = "mhlo.compare"(%5, %0) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<f32>, tensor<f32>) -> tensor<i1> %7 = tensor.extract %6[] : tensor<i1> cf.cond_br %7, ^bb2(%4, %5 : tensor<f32>, tensor<f32>), ^bb6(%4 : tensor<f32>) ^bb2(%8: tensor<f32>, %9: tensor<f32>): %10 = mhlo.add %8, %0 : tensor<f32> %11 = mhlo.remainder %9, %2 : tensor<f32> - %12 = "mhlo.compare"(%11, %3) {comparison_direction = "NE"} : (tensor<f32>, tensor<f32>) -> tensor<i1> - %13 = "mhlo.compare"(%2, %3) {comparison_direction = "LT"} : (tensor<f32>, tensor<f32>) -> tensor<i1> - %14 = "mhlo.compare"(%11, %3) {comparison_direction = "LT"} : (tensor<f32>, tensor<f32>) -> tensor<i1> - %15 = "mhlo.compare"(%13, %14) {comparison_direction = "NE"} : (tensor<i1>, tensor<i1>) -> tensor<i1> + %12 = "mhlo.compare"(%11, %3) {comparison_direction = #mhlo<"comparison_direction NE">} : (tensor<f32>, tensor<f32>) -> tensor<i1> + %13 = "mhlo.compare"(%2, %3) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<f32>, tensor<f32>) -> tensor<i1> + %14 = "mhlo.compare"(%11, %3) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<f32>, tensor<f32>) -> tensor<i1> + %15 = "mhlo.compare"(%13, %14) {comparison_direction = #mhlo<"comparison_direction NE">} : (tensor<i1>, tensor<i1>) -> tensor<i1> %16 = mhlo.and %12, %15 : tensor<i1> %17 = mhlo.add %11, %2 : tensor<f32> %18 = "mhlo.select"(%16, %17, %11) : (tensor<i1>, tensor<f32>, tensor<f32>) -> tensor<f32> - %19 = "mhlo.compare"(%18, %3) {comparison_direction = "GT"} : (tensor<f32>, tensor<f32>) -> tensor<i1> + %19 = "mhlo.compare"(%18, %3) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<f32>, tensor<f32>) -> tensor<i1> %20 = tensor.extract %19[] : tensor<i1> cf.cond_br %20, ^bb3, ^bb4 ^bb3: // pred: ^bb2
diff --git a/iree/test/e2e/models/fragment_000.mlir b/iree/test/e2e/models/fragment_000.mlir index b39037a..f941e5f 100644 --- a/iree/test/e2e/models/fragment_000.mlir +++ b/iree/test/e2e/models/fragment_000.mlir
@@ -13,7 +13,7 @@ %6 = "mhlo.broadcast_in_dim"(%2) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.9"} : (tensor<f32>) -> tensor<5x1x5xf32> %7 = mhlo.multiply %5, %6 : tensor<5x1x5xf32> %8 = "mhlo.broadcast_in_dim"(%0) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.47"} : (tensor<f32>) -> tensor<5x1x5xf32> - %9 = "mhlo.compare"(%7, %8) {comparison_direction = "GT"} : (tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xi1> + %9 = "mhlo.compare"(%7, %8) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xi1> %10 = "mhlo.broadcast_in_dim"(%0) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.11"} : (tensor<f32>) -> tensor<5x1x5xf32> %11 = "mhlo.broadcast_in_dim"(%0) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.67"} : (tensor<f32>) -> tensor<5x5xf32> %12 = "mhlo.broadcast_in_dim"(%4) {broadcast_dimensions = dense<1> : tensor<1xi64>, name = "broadcast.64"} : (tensor<5xf32>) -> tensor<5x5xf32>
diff --git a/iree/test/e2e/models/fullyconnected.mlir b/iree/test/e2e/models/fullyconnected.mlir index 36b1376..b3ba4c8 100644 --- a/iree/test/e2e/models/fullyconnected.mlir +++ b/iree/test/e2e/models/fullyconnected.mlir
@@ -13,7 +13,7 @@ %6 = mhlo.multiply %4, %5 {name = "multiply.57"} : tensor<5x1x5xf32> %cst_0 = arith.constant {name = "constant.58"} dense<0.000000e+00> : tensor<f32> %7 = "mhlo.broadcast_in_dim"(%cst_0) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.59"} : (tensor<f32>) -> tensor<5x1x5xf32> - %8 = "mhlo.compare"(%6, %7) {comparison_direction = "GT", name = "compare.60"} : (tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xi1> + %8 = "mhlo.compare"(%6, %7) {comparison_direction = #mhlo<"comparison_direction GT">, name = "compare.60"} : (tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xi1> %cst_1 = arith.constant {name = "constant.24"} dense<0.000000e+00> : tensor<f32> %9 = "mhlo.broadcast_in_dim"(%cst_1) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.25"} : (tensor<f32>) -> tensor<5x1x5xf32> %cst_2 = arith.constant {name = "constant.90"} dense<0.000000e+00> : tensor<f32> @@ -25,7 +25,7 @@ %14 = mhlo.multiply %12, %13 {name = "multiply.51"} : tensor<5x1x5xf32> %cst_4 = arith.constant {name = "constant.52"} dense<0.000000e+00> : tensor<f32> %15 = "mhlo.broadcast_in_dim"(%cst_4) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.53"} : (tensor<f32>) -> tensor<5x1x5xf32> - %16 = "mhlo.compare"(%14, %15) {comparison_direction = "GT", name = "compare.54"} : (tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xi1> + %16 = "mhlo.compare"(%14, %15) {comparison_direction = #mhlo<"comparison_direction GT">, name = "compare.54"} : (tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xi1> %cst_5 = arith.constant {name = "constant.17"} dense<0.000000e+00> : tensor<f32> %17 = "mhlo.broadcast_in_dim"(%cst_5) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.18"} : (tensor<f32>) -> tensor<5x1x5xf32> %cst_6 = arith.constant {name = "constant.78"} dense<0.000000e+00> : tensor<f32> @@ -37,7 +37,7 @@ %22 = mhlo.multiply %20, %21 {name = "multiply.45"} : tensor<5x1x5xf32> %cst_8 = arith.constant {name = "constant.46"} dense<0.000000e+00> : tensor<f32> %23 = "mhlo.broadcast_in_dim"(%cst_8) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.47"} : (tensor<f32>) -> tensor<5x1x5xf32> - %24 = "mhlo.compare"(%22, %23) {comparison_direction = "GT", name = "compare.48"} : (tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xi1> + %24 = "mhlo.compare"(%22, %23) {comparison_direction = #mhlo<"comparison_direction GT">, name = "compare.48"} : (tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xi1> %cst_9 = arith.constant {name = "constant.10"} dense<0.000000e+00> : tensor<f32> %25 = "mhlo.broadcast_in_dim"(%cst_9) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.11"} : (tensor<f32>) -> tensor<5x1x5xf32> %cst_10 = arith.constant {name = "constant.66"} dense<0.000000e+00> : tensor<f32> @@ -48,7 +48,7 @@ %30 = "mhlo.transpose"(%29) {name = "transpose.39", permutation = dense<[1, 0, 2]> : tensor<3xi64>} : (tensor<1x5x3xf32>) -> tensor<5x1x3xf32> %31 = "mhlo.reshape"(%30) {name = "reshape.40"} : (tensor<5x1x3xf32>) -> tensor<5x3xf32> %cst_11 = arith.constant {name = "constant.61"} dense<[[0.706495285, -0.567672312, 0.483717591, 0.522725761, 0.7563259], [-0.0899272263, -0.283501834, -0.350822538, -0.351515919, -0.337136656], [-0.451804549, 0.372324884, -0.620518147, 0.235451385, 0.851095855]]> : tensor<3x5xf32> - %32 = "mhlo.dot"(%31, %cst_11) {name = "dot.62", precision_config = ["DEFAULT", "DEFAULT"]} : (tensor<5x3xf32>, tensor<3x5xf32>) -> tensor<5x5xf32> + %32 = "mhlo.dot"(%31, %cst_11) {name = "dot.62", precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">]} : (tensor<5x3xf32>, tensor<3x5xf32>) -> tensor<5x5xf32> %cst_12 = arith.constant {name = "constant.63"} dense<[0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00]> : tensor<5xf32> %33 = "mhlo.broadcast_in_dim"(%cst_12) {broadcast_dimensions = dense<[1]> : tensor<1xi64>, name = "broadcast.64"} : (tensor<5xf32>) -> tensor<5x5xf32> %34 = mhlo.add %32, %33 {name = "add.65"} : tensor<5x5xf32> @@ -58,7 +58,7 @@ %38 = "mhlo.copy"(%37) {name = "copy.4"} : (tensor<5x1x5xf32>) -> tensor<5x1x5xf32> %39 = "mhlo.reshape"(%38) {name = "reshape.72"} : (tensor<5x1x5xf32>) -> tensor<5x5xf32> %cst_13 = arith.constant {name = "constant.73"} dense<[[-0.0118641369, -3.785000e-02, 0.489048243, 0.321015775, -0.702280283], [-0.280262798, -0.724645615, -0.00332254497, 0.392334729, 0.619746447], [-0.113318317, -0.180415511, -0.146743968, 0.250408649, -0.442881733], [0.115600757, 0.703136146, -0.00812680274, -0.225454301, -0.0835619792], [-0.136745885, -6.298570e-01, 0.43629986, -0.689790308, 0.230725273]]> : tensor<5x5xf32> - %40 = "mhlo.dot"(%39, %cst_13) {name = "dot.74", precision_config = ["DEFAULT", "DEFAULT"]} : (tensor<5x5xf32>, tensor<5x5xf32>) -> tensor<5x5xf32> + %40 = "mhlo.dot"(%39, %cst_13) {name = "dot.74", precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">]} : (tensor<5x5xf32>, tensor<5x5xf32>) -> tensor<5x5xf32> %cst_14 = arith.constant {name = "constant.75"} dense<[0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00]> : tensor<5xf32> %41 = "mhlo.broadcast_in_dim"(%cst_14) {broadcast_dimensions = dense<[1]> : tensor<1xi64>, name = "broadcast.76"} : (tensor<5xf32>) -> tensor<5x5xf32> %42 = mhlo.add %40, %41 {name = "add.77"} : tensor<5x5xf32> @@ -68,7 +68,7 @@ %46 = "mhlo.copy"(%45) {name = "copy.5"} : (tensor<5x1x5xf32>) -> tensor<5x1x5xf32> %47 = "mhlo.reshape"(%46) {name = "reshape.84"} : (tensor<5x1x5xf32>) -> tensor<5x5xf32> %cst_15 = arith.constant {name = "constant.85"} dense<[[-0.136191264, -0.0401721969, 0.38497138, -5.850760e-01, 0.370910525], [-0.391011149, 0.0266356133, 0.309115469, -0.205079094, -0.559861302], [0.497760415, 0.689488232, 0.0759292394, -0.33134672, -0.237128958], [-0.53243047, 0.476418108, -0.371978909, 0.283265263, 0.63842845], [0.101761498, -0.218626946, 0.475128263, 0.042601984, 0.0988005772]]> : tensor<5x5xf32> - %48 = "mhlo.dot"(%47, %cst_15) {name = "dot.86", precision_config = ["DEFAULT", "DEFAULT"]} : (tensor<5x5xf32>, tensor<5x5xf32>) -> tensor<5x5xf32> + %48 = "mhlo.dot"(%47, %cst_15) {name = "dot.86", precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">]} : (tensor<5x5xf32>, tensor<5x5xf32>) -> tensor<5x5xf32> %cst_16 = arith.constant {name = "constant.87"} dense<[0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00]> : tensor<5xf32> %49 = "mhlo.broadcast_in_dim"(%cst_16) {broadcast_dimensions = dense<[1]> : tensor<1xi64>, name = "broadcast.88"} : (tensor<5xf32>) -> tensor<5x5xf32> %50 = mhlo.add %48, %49 {name = "add.89"} : tensor<5x5xf32>
diff --git a/iree/test/e2e/models/unidirectional_lstm.mlir b/iree/test/e2e/models/unidirectional_lstm.mlir index 45ac605..bae3859 100644 --- a/iree/test/e2e/models/unidirectional_lstm.mlir +++ b/iree/test/e2e/models/unidirectional_lstm.mlir
@@ -13,7 +13,7 @@ // some calls from @main and the call graphs of the removed callees. func private @ForwardLoopCond_gFAnjWGSoLs__.167(%arg0: tensor<i64>, %arg1: tensor<i64>, %arg2: tensor<40xf32>, %arg3: tensor<i64>, %arg4: tensor<74x40xf32>, %arg5: tensor<i64>, %arg6: tensor<1x10xf32>, %arg7: tensor<1x10xf32>, %arg8: tensor<5x1x64xf32>, %arg9: tensor<5x1x1xf32>, %arg10: tensor<5x1x1xf32>, %arg11: tensor<5xi64>, %arg12: tensor<5x1x10xf32>, %arg13: tensor<5x1x10xf32>) -> tensor<i1> { - %0 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<i64>, tensor<i64>) -> tensor<i1> + %0 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<i64>, tensor<i64>) -> tensor<i1> return %0 : tensor<i1> } func private @Forward_o16DF3vQKaI__disable_call_shape_inference_true_.189(%arg0: tensor<1x10xf32>, %arg1: tensor<1x10xf32>, %arg2: tensor<5x1x64xf32>, %arg3: tensor<5x1x1xf32>, %arg4: tensor<5x1x1xf32>) -> (tensor<i64>, tensor<5xi64>, tensor<5x1x10xf32>, tensor<5x1x10xf32>, tensor<i64>, tensor<1x10xf32>, tensor<1x10xf32>) { @@ -39,7 +39,7 @@ %115 = mhlo.minimum %arg5, %arg6 : tensor<f32> "mhlo.return"(%115) : (tensor<f32>) -> () }) {dimensions = dense<[1, 2]> : tensor<2xi64>} : (tensor<5x1x1xf32>, tensor<f32>) -> tensor<5xf32> - %10 = "mhlo.compare"(%9, %0) {comparison_direction = "EQ"} : (tensor<5xf32>, tensor<5xf32>) -> tensor<5xi1> + %10 = "mhlo.compare"(%9, %0) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<5xf32>, tensor<5xf32>) -> tensor<5xi1> %11 = "mhlo.convert"(%10) : (tensor<5xi1>) -> tensor<5xi32> %12 = mhlo.multiply %11, %cst_0 : tensor<5xi32> %13 = "mhlo.reduce"(%12, %cst_1) ( { @@ -48,9 +48,9 @@ "mhlo.return"(%115) : (tensor<i32>) -> () }) {dimensions = dense<0> : tensor<1xi64>} : (tensor<5xi32>, tensor<i32>) -> tensor<i32> %14 = mhlo.subtract %cst_2, %13 : tensor<i32> - %15 = "mhlo.compare"(%14, %cst_2) {comparison_direction = "EQ"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %15 = "mhlo.compare"(%14, %cst_2) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i32>, tensor<i32>) -> tensor<i1> %16 = "mhlo.reverse"(%9) {dimensions = dense<0> : tensor<1xi64>} : (tensor<5xf32>) -> tensor<5xf32> - %17 = "mhlo.compare"(%16, %0) {comparison_direction = "EQ"} : (tensor<5xf32>, tensor<5xf32>) -> tensor<5xi1> + %17 = "mhlo.compare"(%16, %0) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<5xf32>, tensor<5xf32>) -> tensor<5xi1> %18 = "mhlo.convert"(%17) : (tensor<5xi1>) -> tensor<5xi32> %19 = mhlo.multiply %18, %cst_0 : tensor<5xi32> %20 = "mhlo.reduce"(%19, %cst_1) ( { @@ -82,7 +82,7 @@ %58 = "mhlo.reshape"(%57) : (tensor<1x1xf32>) -> tensor<1xf32> %59 = "mhlo.broadcast_in_dim"(%58) {broadcast_dimensions = dense<0> : tensor<1xi64>} : (tensor<1xf32>) -> tensor<1x10xf32> %60 = mhlo.multiply %59, %6 : tensor<1x10xf32> - %61 = "mhlo.compare"(%60, %7) {comparison_direction = "GT"} : (tensor<1x10xf32>, tensor<1x10xf32>) -> tensor<1x10xi1> + %61 = "mhlo.compare"(%60, %7) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<1x10xf32>, tensor<1x10xf32>) -> tensor<1x10xi1> %62 = "mhlo.gather"(%50, %42) { dimension_numbers = #mhlo.gather< collapsed_slice_dims = [0], @@ -93,7 +93,7 @@ slice_sizes = dense<[1, 1, 64]> : tensor<3xi64> } : (tensor<5x1x64xf32>, tensor<i64>) -> tensor<1x64xf32> %63 = "mhlo.concatenate"(%62, %49) {dimension = 1 : i64} : (tensor<1x64xf32>, tensor<1x10xf32>) -> tensor<1x74xf32> - %64 = "mhlo.dot"(%63, %46) {precision_config = ["DEFAULT", "DEFAULT"]} : (tensor<1x74xf32>, tensor<74x40xf32>) -> tensor<1x40xf32> + %64 = "mhlo.dot"(%63, %46) {precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">]} : (tensor<1x74xf32>, tensor<74x40xf32>) -> tensor<1x40xf32> %65 = "mhlo.reshape"(%44) : (tensor<40xf32>) -> tensor<1x40xf32> %66 = mhlo.add %64, %65 : tensor<1x40xf32> %67 = "mhlo.slice"(%66) {limit_indices = dense<[1, 30]> : tensor<2xi64>, start_indices = dense<[0, 20]> : tensor<2xi64>, strides = dense<1> : tensor<2xi64>} : (tensor<1x40xf32>) -> tensor<1x10xf32> @@ -117,7 +117,7 @@ %85 = "mhlo.reshape"(%57) : (tensor<1x1xf32>) -> tensor<1xf32> %86 = "mhlo.broadcast_in_dim"(%85) {broadcast_dimensions = dense<0> : tensor<1xi64>} : (tensor<1xf32>) -> tensor<1x10xf32> %87 = mhlo.multiply %86, %6 : tensor<1x10xf32> - %88 = "mhlo.compare"(%87, %7) {comparison_direction = "GT"} : (tensor<1x10xf32>, tensor<1x10xf32>) -> tensor<1x10xi1> + %88 = "mhlo.compare"(%87, %7) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<1x10xf32>, tensor<1x10xf32>) -> tensor<1x10xi1> %89 = "mhlo.slice"(%66) {limit_indices = dense<[1, 40]> : tensor<2xi64>, start_indices = dense<[0, 30]> : tensor<2xi64>, strides = dense<1> : tensor<2xi64>} : (tensor<1x40xf32>) -> tensor<1x10xf32> %90 = mhlo.multiply %89, %8 : tensor<1x10xf32> %91 = "mhlo.tanh"(%90) : (tensor<1x10xf32>) -> tensor<1x10xf32>
diff --git a/iree/test/e2e/regression/dynamic_compare_and_select.mlir b/iree/test/e2e/regression/dynamic_compare_and_select.mlir index f1ac3a5..01a483b 100644 --- a/iree/test/e2e/regression/dynamic_compare_and_select.mlir +++ b/iree/test/e2e/regression/dynamic_compare_and_select.mlir
@@ -6,7 +6,7 @@ // CHECK: 10xi32=9 8 7 6 5 4 3 2 1 0 func @main(%arg0: tensor<?xi32>, %arg1: tensor<?xi32>, %arg2: tensor<?xi32>, %arg3: tensor<?xi32>) -> tensor<?xi32> { - %1 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<?xi32>, tensor<?xi32>) -> tensor<?xi1> + %1 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<?xi32>, tensor<?xi32>) -> tensor<?xi1> %2 = "mhlo.select"(%1, %arg2, %arg3) : (tensor<?xi1>, tensor<?xi32>, tensor<?xi32>) -> tensor<?xi32> return %2 : tensor<?xi32> }
diff --git a/iree/test/e2e/vulkan_specific/compare.mlir b/iree/test/e2e/vulkan_specific/compare.mlir index 0f171be..12c45e5 100644 --- a/iree/test/e2e/vulkan_specific/compare.mlir +++ b/iree/test/e2e/vulkan_specific/compare.mlir
@@ -1,7 +1,7 @@ func @compare_tensor() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -12,7 +12,7 @@ func @compare_scalar() { %lhs = util.unfoldable_constant dense<1> : tensor<i32> %rhs = util.unfoldable_constant dense<5> : tensor<i32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i32>, tensor<i32>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -23,7 +23,7 @@ func @compare_i8() { %lhs = util.unfoldable_constant dense<1> : tensor<i8> %rhs = util.unfoldable_constant dense<5> : tensor<i8> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<i8>, tensor<i8>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i8>, tensor<i8>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -34,7 +34,7 @@ func @compare_i16() { %lhs = util.unfoldable_constant dense<1> : tensor<i16> %rhs = util.unfoldable_constant dense<5> : tensor<i16> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<i16>, tensor<i16>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i16>, tensor<i16>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -45,7 +45,7 @@ func @compare_i32() { %lhs = util.unfoldable_constant dense<1> : tensor<i32> %rhs = util.unfoldable_constant dense<5> : tensor<i32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i32>, tensor<i32>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -56,7 +56,7 @@ func @compare_i64() { %lhs = util.unfoldable_constant dense<1> : tensor<i64> %rhs = util.unfoldable_constant dense<5> : tensor<i64> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<i64>, tensor<i64>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i64>, tensor<i64>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -67,7 +67,7 @@ func @compare_f32() { %lhs = util.unfoldable_constant dense<1.0> : tensor<f32> %rhs = util.unfoldable_constant dense<5.0> : tensor<f32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<f32>, tensor<f32>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<f32>, tensor<f32>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -78,7 +78,7 @@ func @compare_f64() { %lhs = util.unfoldable_constant dense<1.0> : tensor<f64> %rhs = util.unfoldable_constant dense<5.0> : tensor<f64> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<f64>, tensor<f64>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<f64>, tensor<f64>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -89,7 +89,7 @@ func @compare_tensor_odd_length() { %lhs = util.unfoldable_constant dense<[1, 2, 7]> : tensor<3xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3]> : tensor<3xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<3xi32>, tensor<3xi32>) -> tensor<3xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<3xi32>, tensor<3xi32>) -> tensor<3xi1> %c0 = util.unfoldable_constant dense<0> : tensor<3xi8> %c1 = util.unfoldable_constant dense<1> : tensor<3xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<3xi1>, tensor<3xi8>, tensor<3xi8>) -> tensor<3xi8> @@ -100,7 +100,7 @@ func @compare_eq() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -111,7 +111,7 @@ func @compare_ne() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "NE"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction NE">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -122,7 +122,7 @@ func @compare_lt() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "LT"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -133,7 +133,7 @@ func @compare_le() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "LE"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction LE">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -144,7 +144,7 @@ func @compare_gt() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "GT"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -155,7 +155,7 @@ func @compare_ge() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "GE"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction GE">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8>
diff --git a/iree/test/e2e/vulkan_specific/dot_general.mlir b/iree/test/e2e/vulkan_specific/dot_general.mlir index af2b39c..060af86 100644 --- a/iree/test/e2e/vulkan_specific/dot_general.mlir +++ b/iree/test/e2e/vulkan_specific/dot_general.mlir
@@ -11,7 +11,7 @@ rhs_batching_dimensions = [0], rhs_contracting_dimensions = [1], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<1x2x3xf32>, tensor<1x3x4xf32>) -> tensor<1x2x4xf32> check.expect_almost_eq_const(%res, dense<[[[0.6, 1.2, 1.8, 2.4],[1.5, 3.0, 4.5, 6.0]]]> : tensor<1x2x4xf32>) : tensor<1x2x4xf32> return @@ -36,7 +36,7 @@ rhs_batching_dimensions = [0], rhs_contracting_dimensions = [1], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<2x2x3xf32>, tensor<2x3x4xf32>) -> tensor<2x2x4xf32> check.expect_almost_eq_const(%res, dense<[ [ @@ -58,7 +58,7 @@ rhs_batching_dimensions = [0], rhs_contracting_dimensions = [1], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<4x32x1024xf32>, tensor<4x1024x64xf32>) -> tensor<4x32x64xf32> check.expect_almost_eq_const(%res, dense<409.596> : tensor<4x32x64xf32>) : tensor<4x32x64xf32> return
diff --git a/iree/test/e2e/xla_ops/compare.mlir b/iree/test/e2e/xla_ops/compare.mlir index 0f171be..12c45e5 100644 --- a/iree/test/e2e/xla_ops/compare.mlir +++ b/iree/test/e2e/xla_ops/compare.mlir
@@ -1,7 +1,7 @@ func @compare_tensor() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -12,7 +12,7 @@ func @compare_scalar() { %lhs = util.unfoldable_constant dense<1> : tensor<i32> %rhs = util.unfoldable_constant dense<5> : tensor<i32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i32>, tensor<i32>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -23,7 +23,7 @@ func @compare_i8() { %lhs = util.unfoldable_constant dense<1> : tensor<i8> %rhs = util.unfoldable_constant dense<5> : tensor<i8> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<i8>, tensor<i8>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i8>, tensor<i8>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -34,7 +34,7 @@ func @compare_i16() { %lhs = util.unfoldable_constant dense<1> : tensor<i16> %rhs = util.unfoldable_constant dense<5> : tensor<i16> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<i16>, tensor<i16>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i16>, tensor<i16>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -45,7 +45,7 @@ func @compare_i32() { %lhs = util.unfoldable_constant dense<1> : tensor<i32> %rhs = util.unfoldable_constant dense<5> : tensor<i32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i32>, tensor<i32>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -56,7 +56,7 @@ func @compare_i64() { %lhs = util.unfoldable_constant dense<1> : tensor<i64> %rhs = util.unfoldable_constant dense<5> : tensor<i64> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<i64>, tensor<i64>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i64>, tensor<i64>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -67,7 +67,7 @@ func @compare_f32() { %lhs = util.unfoldable_constant dense<1.0> : tensor<f32> %rhs = util.unfoldable_constant dense<5.0> : tensor<f32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<f32>, tensor<f32>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<f32>, tensor<f32>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -78,7 +78,7 @@ func @compare_f64() { %lhs = util.unfoldable_constant dense<1.0> : tensor<f64> %rhs = util.unfoldable_constant dense<5.0> : tensor<f64> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<f64>, tensor<f64>) -> tensor<i1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<f64>, tensor<f64>) -> tensor<i1> %c0 = util.unfoldable_constant dense<0> : tensor<i8> %c1 = util.unfoldable_constant dense<1> : tensor<i8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<i1>, tensor<i8>, tensor<i8>) -> tensor<i8> @@ -89,7 +89,7 @@ func @compare_tensor_odd_length() { %lhs = util.unfoldable_constant dense<[1, 2, 7]> : tensor<3xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3]> : tensor<3xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<3xi32>, tensor<3xi32>) -> tensor<3xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<3xi32>, tensor<3xi32>) -> tensor<3xi1> %c0 = util.unfoldable_constant dense<0> : tensor<3xi8> %c1 = util.unfoldable_constant dense<1> : tensor<3xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<3xi1>, tensor<3xi8>, tensor<3xi8>) -> tensor<3xi8> @@ -100,7 +100,7 @@ func @compare_eq() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "EQ"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -111,7 +111,7 @@ func @compare_ne() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "NE"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction NE">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -122,7 +122,7 @@ func @compare_lt() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "LT"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -133,7 +133,7 @@ func @compare_le() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "LE"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction LE">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -144,7 +144,7 @@ func @compare_gt() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "GT"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8> @@ -155,7 +155,7 @@ func @compare_ge() { %lhs = util.unfoldable_constant dense<[1, 2, 7, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 2, 3, 4]> : tensor<4xi32> - %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = "GE"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %result = "mhlo.compare"(%lhs, %rhs) {comparison_direction = #mhlo<"comparison_direction GE">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %c0 = util.unfoldable_constant dense<0> : tensor<4xi8> %c1 = util.unfoldable_constant dense<1> : tensor<4xi8> %output = "mhlo.select"(%result, %c1, %c0) : (tensor<4xi1>, tensor<4xi8>, tensor<4xi8>) -> tensor<4xi8>
diff --git a/iree/test/e2e/xla_ops/dot_general.mlir b/iree/test/e2e/xla_ops/dot_general.mlir index d56b98a..297aabb 100644 --- a/iree/test/e2e/xla_ops/dot_general.mlir +++ b/iree/test/e2e/xla_ops/dot_general.mlir
@@ -8,7 +8,7 @@ rhs_batching_dimensions = [], rhs_contracting_dimensions = [0], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<1x1x2xf32>, tensor<2x3xf32>) -> tensor<1x1x3xf32> check.expect_almost_eq_const(%res, dense<[[[0.23, 0.31, 0.39]]]> : tensor<1x1x3xf32>) : tensor<1x1x3xf32> return @@ -24,7 +24,7 @@ rhs_batching_dimensions = [], rhs_contracting_dimensions = [2], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<2x3xf32>, tensor<1x1x2xf32>) -> tensor<3x1x1xf32> check.expect_almost_eq_const(%res, dense<[[[0.23]],[[0.31]],[[0.39]]]> : tensor<3x1x1xf32>) : tensor<3x1x1xf32> return @@ -43,7 +43,7 @@ rhs_batching_dimensions = [0], rhs_contracting_dimensions = [1], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<1x2x3xf32>, tensor<1x3x4xf32>) -> tensor<1x2x4xf32> check.expect_almost_eq_const(%res, dense<[[[0.6, 1.2, 1.8, 2.4],[1.5, 3.0, 4.5, 6.0]]]> : tensor<1x2x4xf32>) : tensor<1x2x4xf32> return @@ -98,7 +98,7 @@ rhs_batching_dimensions = [0], rhs_contracting_dimensions = [1], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<2x2x3xf32>, tensor<2x3x4xf32>) -> tensor<2x2x4xf32> check.expect_almost_eq_const(%res, dense<[ [ @@ -120,7 +120,7 @@ rhs_batching_dimensions = [0], rhs_contracting_dimensions = [1], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<4x32x1024xf32>, tensor<4x1024x64xf32>) -> tensor<4x32x64xf32> check.expect_almost_eq_const(%res, dense<409.596> : tensor<4x32x64xf32>) : tensor<4x32x64xf32> return @@ -136,7 +136,7 @@ rhs_batching_dimensions = [0], rhs_contracting_dimensions = [1], >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<4x32x1024xf32>, tensor<4x1024x64xf32>) -> tensor<4x32x64xf32> check.expect_almost_eq_const(%res, dense<409.596> : tensor<4x32x64xf32>) : tensor<4x32x64xf32> return @@ -157,7 +157,7 @@ lhs_contracting_dimensions = [3], rhs_contracting_dimensions = [2] >, - precision_config = ["DEFAULT", "DEFAULT"] + precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">] } : (tensor<2x3x2x1xf32>, tensor<2x2x1xf32>) -> tensor<2x2x3x2xf32> check.expect_almost_eq_const(%res, dense<[ [
diff --git a/iree/test/e2e/xla_ops/fft.mlir b/iree/test/e2e/xla_ops/fft.mlir index 92d690a..01fc2b6 100644 --- a/iree/test/e2e/xla_ops/fft.mlir +++ b/iree/test/e2e/xla_ops/fft.mlir
@@ -7,7 +7,7 @@ // 3.5, -4.5, 0.0, 9.0, 1.0, 4.5, -0.3, 10.0, -1.0, 5.5, 0.3, 299.0, 3.5, // -0.777, 2.0, 1.7, 3.5, -4.5, 0.0]> : tensor<32xf32> // %0 = "mhlo.fft"(%input) { -// fft_length = dense<32> : tensor<i64>, fft_type = "RFFT" +// fft_length = dense<32> : tensor<i64>, fft_type = #mhlo<"fft_type RFFT"> // } : (tensor<32xf32>) -> tensor<17xcomplex<f32>> // %1 = "mhlo.real"(%0) : (tensor<17xcomplex<f32>>) -> tensor<17xf32> // %2 = "mhlo.imag"(%0) : (tensor<17xcomplex<f32>>) -> tensor<17xf32> @@ -22,7 +22,7 @@ 3.5, -4.5, 0.0, 9.0, 1.0, 4.5, -0.3, 10.0, -1.0, 5.5, 0.3, 299.0, 3.5, -0.777, 2.0, 1.7, 3.5, -4.5, 0.0]]> : tensor<1x32xf32> %0 = "mhlo.fft"(%input) { - fft_length = dense<32> : tensor<1xi64>, fft_type = "RFFT" + fft_length = dense<32> : tensor<1xi64>, fft_type = #mhlo<"fft_type RFFT"> } : (tensor<1x32xf32>) -> tensor<1x17xcomplex<f32>> %1 = "mhlo.real"(%0) : (tensor<1x17xcomplex<f32>>) -> tensor<1x17xf32> %2 = "mhlo.imag"(%0) : (tensor<1x17xcomplex<f32>>) -> tensor<1x17xf32>
diff --git a/iree/test/e2e/xla_ops/reduce.mlir b/iree/test/e2e/xla_ops/reduce.mlir index f0b9e53..891ad93 100644 --- a/iree/test/e2e/xla_ops/reduce.mlir +++ b/iree/test/e2e/xla_ops/reduce.mlir
@@ -296,9 +296,9 @@ %arg1 = util.unfoldable_constant dense<[[0, 1], [2, 3], [4, 5], [6, 7], [8, 9], [10, 11], [12, 13], [14, 15], [16, 17]]> : tensor<9x2xi32> %res0, %res1 = "mhlo.reduce"(%arg0, %arg1, %cst0, %cst1) ( { ^bb0(%arg2: tensor<i32>, %arg3: tensor<i32>, %arg4: tensor<i32>, %arg5: tensor<i32>): // no predecessors - %0 = "mhlo.compare"(%arg2, %arg4) {comparison_direction = "GE"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %0 = "mhlo.compare"(%arg2, %arg4) {comparison_direction = #mhlo<"comparison_direction GE">} : (tensor<i32>, tensor<i32>) -> tensor<i1> %1 = "mhlo.select"(%0, %arg2, %arg4) : (tensor<i1>, tensor<i32>, tensor<i32>) -> tensor<i32> - %2 = "mhlo.compare"(%arg2, %arg4) {comparison_direction = "EQ"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %2 = "mhlo.compare"(%arg2, %arg4) {comparison_direction = #mhlo<"comparison_direction EQ">} : (tensor<i32>, tensor<i32>) -> tensor<i1> %3 = mhlo.minimum %arg3, %arg5 : tensor<i32> %4 = "mhlo.select"(%0, %arg3, %arg5) : (tensor<i1>, tensor<i32>, tensor<i32>) -> tensor<i32> %5 = "mhlo.select"(%2, %3, %4) : (tensor<i1>, tensor<i32>, tensor<i32>) -> tensor<i32>
diff --git a/iree/test/e2e/xla_ops/select.mlir b/iree/test/e2e/xla_ops/select.mlir index d2cff05..2819987 100644 --- a/iree/test/e2e/xla_ops/select.mlir +++ b/iree/test/e2e/xla_ops/select.mlir
@@ -2,7 +2,7 @@ // TODO(b/132205704) support i1 in constants and function signatures. %input = util.unfoldable_constant dense<[1, 0, 1, 0]> : tensor<4xi32> %zeros = util.unfoldable_constant dense<0> : tensor<4xi32> - %cond = "mhlo.compare"(%input, %zeros) {comparison_direction = "GT"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> + %cond = "mhlo.compare"(%input, %zeros) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> %lhs = util.unfoldable_constant dense<[1, 2, 3, 4]> : tensor<4xi32> %rhs = util.unfoldable_constant dense<[5, 6, 7, 8]> : tensor<4xi32> %result = "mhlo.select"(%cond, %lhs, %rhs) : (tensor<4xi1>, tensor<4xi32>, tensor<4xi32>) -> tensor<4xi32>
diff --git a/iree/test/e2e/xla_ops/sort.mlir b/iree/test/e2e/xla_ops/sort.mlir index 58813ea..b265682 100644 --- a/iree/test/e2e/xla_ops/sort.mlir +++ b/iree/test/e2e/xla_ops/sort.mlir
@@ -3,7 +3,7 @@ %sort = "mhlo.sort"(%input) ( { ^bb0(%arg1: tensor<i32>, %arg2: tensor<i32>): // no predecessors - %compare = "mhlo.compare"(%arg1, %arg2) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %compare = "mhlo.compare"(%arg1, %arg2) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> "mhlo.return"(%compare) : (tensor<i1>) -> () }) {dimension = 0 : i64, is_stable = false} : (tensor<4xi32>) -> tensor<4xi32> @@ -17,7 +17,7 @@ %sort = "mhlo.sort"(%input) ( { ^bb0(%arg1: tensor<i32>, %arg2: tensor<i32>): // no predecessors - %compare = "mhlo.compare"(%arg1, %arg2) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %compare = "mhlo.compare"(%arg1, %arg2) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> "mhlo.return"(%compare) : (tensor<i1>) -> () }) {dimension = 1 : i64, is_stable = false} : (tensor<2x4xi32>) -> tensor<2x4xi32> @@ -31,7 +31,7 @@ %sort = "mhlo.sort"(%input) ( { ^bb0(%arg1: tensor<i32>, %arg2: tensor<i32>): // no predecessors - %compare = "mhlo.compare"(%arg1, %arg2) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %compare = "mhlo.compare"(%arg1, %arg2) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> "mhlo.return"(%compare) : (tensor<i1>) -> () }) {dimension = 2 : i64, is_stable = false} : (tensor<1x2x4xi32>) -> tensor<1x2x4xi32> @@ -44,7 +44,7 @@ %sort = "mhlo.sort"(%input) ( { ^bb0(%arg1: tensor<i32>, %arg2: tensor<i32>): // no predecessors - %compare = "mhlo.compare"(%arg1, %arg2) {comparison_direction = "GT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %compare = "mhlo.compare"(%arg1, %arg2) {comparison_direction = #mhlo<"comparison_direction GT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> "mhlo.return"(%compare) : (tensor<i1>) -> () }) {dimension = 0 : i64, is_stable = false} : (tensor<4xi32>) -> tensor<4xi32>
diff --git a/iree/test/e2e/xla_ops/while.mlir b/iree/test/e2e/xla_ops/while.mlir index 00d02e0..480087f 100644 --- a/iree/test/e2e/xla_ops/while.mlir +++ b/iree/test/e2e/xla_ops/while.mlir
@@ -5,7 +5,7 @@ %cst_1 = arith.constant dense<4> : tensor<i32> cf.br ^bb1(%start : tensor<i32>) ^bb1(%2: tensor<i32>): - %3 = "mhlo.compare"(%2, %bound) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> + %3 = "mhlo.compare"(%2, %bound) {comparison_direction = #mhlo<"comparison_direction LT">} : (tensor<i32>, tensor<i32>) -> tensor<i1> %4 = tensor.extract %3[] : tensor<i1> cf.cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>) ^bb2(%5: tensor<i32>):
diff --git a/llvm-external-projects/iree-dialects/BUILD b/llvm-external-projects/iree-dialects/BUILD index 3f47bd6..3280a69 100644 --- a/llvm-external-projects/iree-dialects/BUILD +++ b/llvm-external-projects/iree-dialects/BUILD
@@ -631,6 +631,7 @@ "@llvm-project//mlir:ArithmeticDialect", "@llvm-project//mlir:BufferizationDialect", "@llvm-project//mlir:BufferizationTransforms", + "@llvm-project//mlir:FuncDialect", "@llvm-project//mlir:LinalgOps", "@llvm-project//mlir:LLVMDialect", "@llvm-project//mlir:PDLDialect",
diff --git a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/PassDetail.h b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/PassDetail.h index 3f3fe9b..417a582 100644 --- a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/PassDetail.h +++ b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/PassDetail.h
@@ -7,6 +7,7 @@ #ifndef IREE_DIALECTS_DIALECT_LINALGEXT_TRANSFORMS_PASS_DETAIL_H_ #define IREE_DIALECTS_DIALECT_LINALGEXT_TRANSFORMS_PASS_DETAIL_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir {
diff --git a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.h b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.h index febec87..98e37ac 100644 --- a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.h +++ b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.h
@@ -7,6 +7,7 @@ #ifndef IREE_DIALECTS_DIALECT_LINALGEXT_TRANSFORMS_PASSES_H_ #define IREE_DIALECTS_DIALECT_LINALGEXT_TRANSFORMS_PASSES_H_ +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" namespace mlir { @@ -14,9 +15,9 @@ namespace IREE { namespace LinalgExt { -std::unique_ptr<OperationPass<FuncOp>> createTiledOpInterfaceTilingPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createTiledOpInterfaceTilingPass(); -std::unique_ptr<OperationPass<FuncOp>> createLinalgExtToLoopsPass(); +std::unique_ptr<OperationPass<func::FuncOp>> createLinalgExtToLoopsPass(); std::unique_ptr<OperationPass<>> createPadContractionToBlockSizePass();
diff --git a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.td b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.td index 54a0484..95ae571 100644 --- a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.td +++ b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Passes/Passes.td
@@ -10,13 +10,13 @@ include "mlir/Pass/PassBase.td" def LinalgExtToLoops : - Pass<"iree-linalg-ext-to-loops", "FuncOp"> { + Pass<"iree-linalg-ext-to-loops", "func::FuncOp"> { let summary = "Convert LinalgExt ops to loops and Linalg ops."; let constructor = "mlir::iree_compiler::IREE::LinalgExt::createLinalgExtToLoopsPass()"; } def TiledOpInterfaceTiling : - Pass<"iree-linalg-ext-tile", "FuncOp"> { + Pass<"iree-linalg-ext-tile", "func::FuncOp"> { let summary = "Test pass for tiling using TiledOpInterface"; let constructor = "mlir::iree_compiler::IREE::LinalgExt::createTiledOpInterfaceTilingPass()"; }
diff --git a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgTransform/LinalgTransformOps.h b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgTransform/LinalgTransformOps.h index 6eda492..7683ae0 100644 --- a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgTransform/LinalgTransformOps.h +++ b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgTransform/LinalgTransformOps.h
@@ -11,6 +11,7 @@ #include "iree-dialects/Dialect/LinalgExt/IR/LinalgExtOps.h" #include "iree-dialects/Dialect/LinalgTransform/TrackingListener.h" #include "iree-dialects/Dialect/LinalgTransform/TransformOpInterface.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/PDL/IR/PDLTypes.h" #include "mlir/IR/BuiltinAttributes.h"
diff --git a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.h b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.h index 905f3a3..c038c0b 100644 --- a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.h +++ b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.h
@@ -12,10 +12,12 @@ #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/Dialect.h" +#include "mlir/IR/FunctionInterfaces.h" #include "mlir/IR/OpDefinition.h" #include "mlir/IR/OpImplementation.h" #include "mlir/IR/PatternMatch.h" #include "mlir/IR/SymbolTable.h" +#include "mlir/Interfaces/CallInterfaces.h" #include "mlir/Interfaces/ControlFlowInterfaces.h" #include "mlir/Interfaces/SideEffectInterfaces.h"
diff --git a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.td b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.td index bc5b181..9b4aca3 100644 --- a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.td +++ b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/PyDM/IR/PyDMOps.td
@@ -124,7 +124,7 @@ }]; let arguments = (ins SymbolNameAttr:$sym_name, - TypeAttr:$type, + TypeAttr:$function_type, OptionalAttr<StrArrayAttr>:$arg_names, OptionalAttr<StrArrayAttr>:$free_vars, OptionalAttr<StrArrayAttr>:$cell_vars, @@ -143,21 +143,21 @@ } /// Returns the type of this function. - FunctionType getType() { + FunctionType getFunctionType() { return getOperation()->getAttrOfType<TypeAttr>(getTypeAttrName()) .getValue() .cast<FunctionType>(); } /// Returns the argument types of this function. - ArrayRef<Type> getArgumentTypes() { return getType().getInputs(); } + ArrayRef<Type> getArgumentTypes() { return getFunctionType().getInputs(); } /// Returns the result types of this function. - ArrayRef<Type> getResultTypes() { return getType().getResults(); } + ArrayRef<Type> getResultTypes() { return getFunctionType().getResults(); } /// Returns the python return type of the function (second return type). Type getPyReturnType() { - return getType().getResult(1); + return getFunctionType().getResult(1); } /// Hook for Trait::FunctionLike, called after verifying that the 'type' @@ -167,7 +167,7 @@ Region *getCallableRegion() { return &body(); } ArrayRef<Type> getCallableResults() { - return getType().getResults(); + return getFunctionType().getResults(); } /// Defines SymbolOpInterface::isDeclaration(). @@ -177,8 +177,8 @@ }]; let builders = [ - OpBuilder<(ins "StringAttr":$name, "FunctionType":$type), [{ - build($_builder, $_state, name, TypeAttr::get(type), + OpBuilder<(ins "StringAttr":$name, "FunctionType":$function_type), [{ + build($_builder, $_state, name, TypeAttr::get(function_type), nullptr, nullptr, nullptr, nullptr); }]> ];
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/IR/TiledOpInterface.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/IR/TiledOpInterface.cpp index e509489..604e498 100644 --- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/IR/TiledOpInterface.cpp +++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/IR/TiledOpInterface.cpp
@@ -305,9 +305,12 @@ DialectRegistry ®istry) { LLVM_DEBUG( { llvm::dbgs() << "Adding external models of tiled op interface\n"; }); - registry - .addOpInterface<tensor::ExtractSliceOp, ExtractSliceTiledOpInterface>(); - registry.addOpInterface<tensor::InsertSliceOp, InsertSliceTiledOpInterface>(); + + registry.addExtension(+[](MLIRContext *ctx, tensor::TensorDialect *dialect) { + tensor::ExtractSliceOp::attachInterface<ExtractSliceTiledOpInterface>(*ctx); + tensor::InsertSliceOp::attachInterface<InsertSliceTiledOpInterface>(*ctx); + }); + // TODO(ravishankarm): Needs custom PadTiledOpInterface or equiv. // registry.addOpInterface<tensor::PadOp, // ForwardToTilingInterface<tensor::PadOp>>();
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/LinalgExtBufferization.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/LinalgExtBufferization.cpp index 2b8f8ec..9b3c2fb 100644 --- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/LinalgExtBufferization.cpp +++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/LinalgExtBufferization.cpp
@@ -6,11 +6,11 @@ #include "iree-dialects/Dialect/LinalgExt/LinalgExtBufferization.h" -#include <mlir/IR/BuiltinOps.h> - +#include "iree-dialects/Dialect/LinalgExt/IR/LinalgExtDialect.h" #include "iree-dialects/Dialect/LinalgExt/IR/LinalgExtOps.h" #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" +#include "mlir/IR/BuiltinOps.h" #include "mlir/IR/PatternMatch.h" using namespace mlir; @@ -22,7 +22,6 @@ using bufferization::BufferRelation; using bufferization::getMemRefType; using bufferization::replaceOpWithBufferizedValues; -using bufferization::replaceOpWithNewBufferizedOp; using tensor::ExtractSliceOp; /// Return the destinations that an InParallelOp is inserting into. One per @@ -341,9 +340,12 @@ void mlir::iree_compiler::IREE::LinalgExt:: registerBufferizableOpInterfaceExternalModels(DialectRegistry ®istry) { - registry.addOpInterface<InParallelOp, InParallelOpInterface>(); - registry - .addOpInterface<PerformConcurrentlyOp, PerformConcurrentlyOpInterface>(); - registry - .addOpInterface<ParallelInsertSliceOp, ParallelInsertSliceOpInterface>(); + registry.addExtension( + +[](MLIRContext *ctx, LinalgExt::IREELinalgExtDialect *dialect) { + InParallelOp::attachInterface<InParallelOpInterface>(*ctx); + PerformConcurrentlyOp::attachInterface<PerformConcurrentlyOpInterface>( + *ctx); + ParallelInsertSliceOp::attachInterface<ParallelInsertSliceOpInterface>( + *ctx); + }); }
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/TilingExternalModels.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/TilingExternalModels.cpp index 174d4ff..0758d7f 100644 --- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/TilingExternalModels.cpp +++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/TilingExternalModels.cpp
@@ -157,23 +157,25 @@ } // namespace template <typename OpType> -void registerOne(DialectRegistry ®istry) { - registry.addOpInterface<OpType, LinalgOpTilingInterface<OpType>>(); +void registerOne(MLIRContext *ctx) { + OpType::template attachInterface<LinalgOpTilingInterface<OpType>>(*ctx); } /// Variadic helper function. template <typename... OpTypes> -void registerAll(DialectRegistry ®istry) { +void registerAll(MLIRContext *ctx) { // FIXME: In c++17 this can be simplified by using 'fold expressions'. - (void)std::initializer_list<int>{0, (registerOne<OpTypes>(registry), 0)...}; + (void)std::initializer_list<int>{0, (registerOne<OpTypes>(ctx), 0)...}; } #define GET_OP_LIST void mlir::iree_compiler::IREE::LinalgExt:: registerTilingInterfaceExternalModels(DialectRegistry ®istry) { - registerOne<linalg::GenericOp>(registry); - registerAll< + registry.addExtension(+[](MLIRContext *ctx, linalg::LinalgDialect *dialect) { + registerOne<linalg::GenericOp>(ctx); + registerAll< #include "mlir/Dialect/Linalg/IR/LinalgStructuredOps.cpp.inc" - >(registry); + >(ctx); + }); }
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgTransform/IR/LinalgTransformOps.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgTransform/IR/LinalgTransformOps.cpp index e28dcd3..03d39ac 100644 --- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgTransform/IR/LinalgTransformOps.cpp +++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgTransform/IR/LinalgTransformOps.cpp
@@ -31,6 +31,7 @@ #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" #include "mlir/Dialect/Bufferization/Transforms/OneShotAnalysis.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/Linalg/ComprehensiveBufferize/ModuleBufferization.h" #include "mlir/Dialect/Linalg/Passes.h" @@ -581,7 +582,7 @@ // Perform buffer-level hoistings. state.getTopLevel()->walk( - [&](FuncOp funcOp) { hoistRedundantVectorTransfers(funcOp); }); + [&](func::FuncOp funcOp) { hoistRedundantVectorTransfers(funcOp); }); return success(); } @@ -597,8 +598,8 @@ // the end. Keep module-level for now. PassManager pm(getContext()); - pm.addNestedPass<FuncOp>(createConvertVectorToSCFPass()); - pm.addNestedPass<FuncOp>(createConvertLinalgToLoopsPass()); + pm.addNestedPass<func::FuncOp>(createConvertVectorToSCFPass()); + pm.addNestedPass<func::FuncOp>(createConvertLinalgToLoopsPass()); if (enable_async()) { pm.addPass(createAsyncToAsyncRuntimePass()); pm.addPass(createAsyncRuntimeRefCountingPass()); @@ -618,7 +619,7 @@ .enableAMX(enable_amx()) .enableX86Vector(enable_x86vector()))); // clang-format on - pm.addNestedPass<FuncOp>(createConvertMathToLLVMPass()); + pm.addNestedPass<func::FuncOp>(createConvertMathToLLVMPass()); pm.addPass(createMemRefToLLVMPass()); if (enable_async()) pm.addPass(createConvertAsyncToLLVMPass()); @@ -631,7 +632,9 @@ // FIXME: this is a terrible hack! state.getTopLevel()->walk([](LLVM::LLVMFuncOp funcOp) { for (int64_t i = 0; i < funcOp.getNumArguments(); ++i) { - if (!funcOp.getType().getParamType(i).isa<LLVM::LLVMPointerType>()) + if (!funcOp.getFunctionType() + .getParamType(i) + .isa<LLVM::LLVMPointerType>()) continue; funcOp.setArgAttr(i, "llvm.noalias", UnitAttr::get(funcOp.getContext())); } @@ -760,15 +763,15 @@ return executeRegionOp; } -static FailureOr<FuncOp> outlineLoop(scf::ForOp loop, StringRef funcName, - transform::TransformState &state) { +static FailureOr<func::FuncOp> outlineLoop(scf::ForOp loop, StringRef funcName, + transform::TransformState &state) { PatternRewriterListener rewriter(loop->getContext()); auto &listener = state.getExtension<TrackingListener>(); rewriter.addListener(&listener); Location loc = loop.getLoc(); scf::ExecuteRegionOp exec = outlineInExecuteRegion(rewriter, loop); assert(exec && "failed to produce execute_region"); - FailureOr<FuncOp> outlined = + FailureOr<func::FuncOp> outlined = outlineSingleBlockRegion(rewriter, loc, exec.getRegion(), funcName); if (failed(listener.checkErrorState())) return failure(); @@ -781,7 +784,7 @@ SmallVector<Operation *> resultVector; auto res = applyTransformToEach(state.getPayloadOps(target()), resultVector, - [&](scf::ForOp loop) -> FailureOr<FuncOp> { + [&](scf::ForOp loop) -> FailureOr<func::FuncOp> { return outlineLoop(loop, func_name(), state); }); if (failed(res))
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/IR/PyDMOps.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/IR/PyDMOps.cpp index b281874..f4ad03a 100644 --- a/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/IR/PyDMOps.cpp +++ b/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/IR/PyDMOps.cpp
@@ -567,9 +567,7 @@ } void PyFuncOp::print(OpAsmPrinter &p) { - FunctionType fnType = getType(); - function_interface_impl::printFunctionOp( - p, *this, fnType.getInputs(), /*isVariadic=*/false, fnType.getResults()); + function_interface_impl::printFunctionOp(p, *this, /*isVariadic=*/false); } //===----------------------------------------------------------------------===// @@ -764,7 +762,7 @@ << "' does not reference a valid function"; // Verify that the operand and result types match the callee. - auto fnType = fn.getType(); + auto fnType = fn.getFunctionType(); if (fnType.getNumInputs() != getNumOperands()) return emitOpError("incorrect number of operands for callee");
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/Transforms/Optimize/FixateWeakNumeric.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/Transforms/Optimize/FixateWeakNumeric.cpp index b6e5b51..5d4be33 100644 --- a/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/Transforms/Optimize/FixateWeakNumeric.cpp +++ b/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/Transforms/Optimize/FixateWeakNumeric.cpp
@@ -41,7 +41,7 @@ // Special cases for operations. if (auto funcOp = llvm::dyn_cast<PYDM::FuncOp>(op)) { - FunctionType existingFt = funcOp.getType(); + FunctionType existingFt = funcOp.getFunctionType(); FunctionType newFt = convertFunctionType(existingFt); if (newFt != existingFt) { funcOp.setType(newFt);
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/Transforms/ToIREE/LoweringPatterns.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/Transforms/ToIREE/LoweringPatterns.cpp index 9d92711..2e14e37 100644 --- a/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/Transforms/ToIREE/LoweringPatterns.cpp +++ b/llvm-external-projects/iree-dialects/lib/Dialect/PyDM/Transforms/ToIREE/LoweringPatterns.cpp
@@ -637,7 +637,7 @@ LogicalResult matchAndRewrite(PYDM::FuncOp srcOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - FunctionType srcFuncType = srcOp.getType(); + FunctionType srcFuncType = srcOp.getFunctionType(); TypeConverter::SignatureConversion signatureConversion( srcOp.getNumArguments()); @@ -839,7 +839,7 @@ auto parentFunc = srcOp->getParentOfType<mlir::FuncOp>(); if (!parentFunc) return rewriter.notifyMatchFailure(srcOp, "not contained by a func"); - Type convertedReturnType = parentFunc.getType().getResult(1); + Type convertedReturnType = parentFunc.getFunctionType().getResult(1); // Split the entry block. Block *entryBlock = rewriter.getInsertionBlock();
diff --git a/llvm-external-projects/iree-dialects/python/iree/compiler/dialects/_iree_pydm_ops_ext.py b/llvm-external-projects/iree-dialects/python/iree/compiler/dialects/_iree_pydm_ops_ext.py index 5155972..7fc8653 100644 --- a/llvm-external-projects/iree-dialects/python/iree/compiler/dialects/_iree_pydm_ops_ext.py +++ b/llvm-external-projects/iree-dialects/python/iree/compiler/dialects/_iree_pydm_ops_ext.py
@@ -18,7 +18,7 @@ @property def type(self): - return ir.FunctionType(ir.TypeAttr(self.attributes["type"]).value) + return ir.FunctionType(ir.TypeAttr(self.attributes["function_type"]).value) @property def py_return_type(self) -> ir.Type:
diff --git a/llvm-external-projects/iree-dialects/python/iree/compiler/dialects/iree_pydm/importer/importer.py b/llvm-external-projects/iree-dialects/python/iree/compiler/dialects/iree_pydm/importer/importer.py index 68b1227..fae6591 100644 --- a/llvm-external-projects/iree-dialects/python/iree/compiler/dialects/iree_pydm/importer/importer.py +++ b/llvm-external-projects/iree-dialects/python/iree/compiler/dialects/iree_pydm/importer/importer.py
@@ -85,7 +85,7 @@ context=ic.context) f_op = d.FuncOp( ir.StringAttr.get(symbol), - type=ir.TypeAttr.get(ir_f_type), + function_type=ir.TypeAttr.get(ir_f_type), arg_names=f_arg_names, free_vars=f_var_names, cell_vars=ir.ArrayAttr.get([]),
diff --git a/llvm-external-projects/iree-dialects/test/Transforms/test-listener-cse.mlir b/llvm-external-projects/iree-dialects/test/Transforms/test-listener-cse.mlir index 13d2994..ebfc348 100644 --- a/llvm-external-projects/iree-dialects/test/Transforms/test-listener-cse.mlir +++ b/llvm-external-projects/iree-dialects/test/Transforms/test-listener-cse.mlir
@@ -233,7 +233,7 @@ %0 = arith.constant 1 : i32 // CHECK-NEXT: @nested_func - builtin.func @nested_func() { + func.func @nested_func() { // CHECK-NEXT: arith.constant 1 %foo = arith.constant 1 : i32 "foo.yield"(%foo) : (i32) -> ()
diff --git a/third_party/llvm-project b/third_party/llvm-project index e9c9ee9..4bbcdb1 160000 --- a/third_party/llvm-project +++ b/third_party/llvm-project
@@ -1 +1 @@ -Subproject commit e9c9ee9fe694067ee96643d05d6ac378349386bb +Subproject commit 4bbcdb1bf868b9288329fa501d498761abcaa92c
diff --git a/third_party/mlir-hlo b/third_party/mlir-hlo index 57288f1..467cd37 160000 --- a/third_party/mlir-hlo +++ b/third_party/mlir-hlo
@@ -1 +1 @@ -Subproject commit 57288f12595a2ee0488806672a42da59b1e56e13 +Subproject commit 467cd37703dc0c4195ce6351617ef320bb60e927