Skip to content

Commit

Permalink
Merge remote-tracking branch 'origin/main' into fill2splat2
Browse files Browse the repository at this point in the history
  • Loading branch information
antiagainst committed Jun 12, 2021
2 parents db6d765 + be6d0b6 commit 47be229
Show file tree
Hide file tree
Showing 199 changed files with 462 additions and 1,609 deletions.
3 changes: 1 addition & 2 deletions bindings/python/iree/compiler/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -6,8 +6,7 @@
import iree.compiler

SIMPLE_MUL_ASM = """
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32>
attributes { iree.module.export } {
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> {
%0 = "mhlo.multiply"(%arg0, %arg1) {name = "mul.1"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
return %0 : tensor<4xf32>
}
Expand Down
3 changes: 1 addition & 2 deletions bindings/python/iree/compiler/setup.py.in
Original file line number Diff line number Diff line change
Expand Up @@ -23,8 +23,7 @@ README = r'''
import iree.compiler
SIMPLE_MUL_ASM = """
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32>
attributes { iree.module.export } {
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> {
%0 = "mhlo.multiply"(%arg0, %arg1) {name = "mul.1"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
return %0 : tensor<4xf32>
}
Expand Down
1 change: 0 additions & 1 deletion bindings/python/iree/runtime/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,6 @@ iree_pyext_module(
"unix_version.lds"
DEPS
iree::base
iree::base::signature_parser
iree::base::status
iree::hal
iree::hal::drivers
Expand Down
3 changes: 1 addition & 2 deletions bindings/python/iree/runtime/system_api_test.py
Original file line number Diff line number Diff line change
Expand Up @@ -20,8 +20,7 @@ def create_simple_mul_module():
binary = iree.compiler.compile_str(
"""
module @arithmetic {
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32>
attributes { iree.module.export } {
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> {
%0 = "mhlo.multiply"(%arg0, %arg1) {name = "mul.1"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
return %0 : tensor<4xf32>
}
Expand Down
8 changes: 3 additions & 5 deletions bindings/python/iree/runtime/vm_test.py
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@
def create_add_scalar_module():
binary = iree.compiler.compile_str(
"""
func @add_scalar(%arg0: i32, %arg1: i32) -> i32 attributes { iree.module.export } {
func @add_scalar(%arg0: i32, %arg1: i32) -> i32 {
%0 = addi %arg0, %arg1 : i32
return %0 : i32
}
Expand All @@ -32,8 +32,7 @@ def create_add_scalar_module():
def create_simple_static_mul_module():
binary = iree.compiler.compile_str(
"""
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32>
attributes { iree.module.export } {
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> {
%0 = "mhlo.multiply"(%arg0, %arg1) {name = "mul.1"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
return %0 : tensor<4xf32>
}
Expand All @@ -50,8 +49,7 @@ def create_simple_dynamic_abs_module():
target_backends = iree.compiler.DEFAULT_TESTING_BACKENDS
binary = iree.compiler.compile_str(
"""
func @simple_mul(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32>
attributes { iree.module.export } {
func @simple_mul(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> {
%0 = "mhlo.abs"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32>
return %0 : tensor<?x?xf32>
}
Expand Down
3 changes: 1 addition & 2 deletions bindings/python/tests/compiler_core_test.py
Original file line number Diff line number Diff line change
Expand Up @@ -15,8 +15,7 @@
import iree.compiler

SIMPLE_MUL_ASM = """
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32>
attributes { iree.module.export } {
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> {
%0 = "mhlo.multiply"(%arg0, %arg1) {name = "mul.1"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
return %0 : tensor<4xf32>
}
Expand Down
4 changes: 0 additions & 4 deletions bindings/python/tests/compiler_xla_test.py
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,6 @@ def testImportBinaryPbFile(self):
import_only=True).decode("utf-8")
logging.info("%s", text)
self.assertIn("linalg.generic", text)
self.assertIn("iree.module.export", text)

def testCompileBinaryPbFile(self):
path = os.path.join(os.path.dirname(__file__), "testdata", "xla_sample.pb")
Expand Down Expand Up @@ -96,7 +95,6 @@ def testImportHloTextFile(self):
path, import_only=True, import_format="hlo_text").decode("utf-8")
logging.info("%s", text)
self.assertIn("linalg.generic", text)
self.assertIn("iree.module.export", text)

def testImportHloTextStr(self):
path = os.path.join(os.path.dirname(__file__), "testdata", "xla_sample.hlo")
Expand All @@ -106,7 +104,6 @@ def testImportHloTextStr(self):
content, import_only=True, import_format="hlo_text").decode("utf-8")
logging.info("%s", text)
self.assertIn("linalg.generic", text)
self.assertIn("iree.module.export", text)

def testImportHloTextBytes(self):
path = os.path.join(os.path.dirname(__file__), "testdata", "xla_sample.hlo")
Expand All @@ -116,7 +113,6 @@ def testImportHloTextBytes(self):
content, import_only=True, import_format="hlo_text").decode("utf-8")
logging.info("%s", text)
self.assertIn("linalg.generic", text)
self.assertIn("iree.module.export", text)


if __name__ == "__main__":
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@ func @main(
%input : tensor<2xf32> {iree.identifier = "input"}
) -> (
tensor<2xf32> {iree.identifier = "output"}
) attributes { iree.module.export } {
) {
%result = mhlo.add %input, %input : tensor<2xf32>
return %result : tensor<2xf32>
}
2 changes: 1 addition & 1 deletion bindings/tflite/testdata/add_dynamic.mlir
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@ func @main(
%input : tensor<?xf32> {iree.identifier = "input"}
) -> (
tensor<?xf32> {iree.identifier = "output"}
) attributes { iree.module.export } {
) {
%result = mhlo.add %input, %input : tensor<?xf32>
return %result : tensor<?xf32>
}
2 changes: 1 addition & 1 deletion bindings/tflite/testdata/add_multi.mlir
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@ func @main(
) -> (
tensor<1x8x8x3xf32> {iree.identifier = "x"},
tensor<1x8x8x3xf32> {iree.identifier = "y"}
) attributes { iree.module.export } {
) {
%0 = "mhlo.add"(%arg1, %arg2) : (tensor<1x8x8x3xf32>, tensor<1x8x8x3xf32>) -> tensor<1x8x8x3xf32>
%1 = "mhlo.add"(%arg0, %0) : (tensor<1x8x8x3xf32>, tensor<1x8x8x3xf32>) -> tensor<1x8x8x3xf32>
%2 = "mhlo.add"(%arg3, %0) : (tensor<1x8x8x3xf32>, tensor<1x8x8x3xf32>) -> tensor<1x8x8x3xf32>
Expand Down
2 changes: 1 addition & 1 deletion bindings/tflite/testdata/add_static.mlir
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@ func @main(
%input : tensor<1x8x8x3xf32> {iree.identifier = "input"}
) -> (
tensor<1x8x8x3xf32> {iree.identifier = "output"}
) attributes { iree.module.export } {
) {
%result = mhlo.add %input, %input : tensor<1x8x8x3xf32>
return %result : tensor<1x8x8x3xf32>
}
4 changes: 2 additions & 2 deletions colab/low_level_invoke_function.ipynb
Original file line number Diff line number Diff line change
Expand Up @@ -111,7 +111,7 @@
"SIMPLE_MUL_ASM = \"\"\"\n",
" module @arithmetic {\n",
" func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32>\n",
" attributes { iree.module.export } {\n",
" {\n",
" %0 = mulf %arg0, %arg1 : tensor<4xf32>\n",
" return %0 : tensor<4xf32>\n",
" } \n",
Expand Down Expand Up @@ -169,4 +169,4 @@
]
}
]
}
}
3 changes: 1 addition & 2 deletions docs/developers/design_docs/cuda_backend.md
Original file line number Diff line number Diff line change
Expand Up @@ -68,8 +68,7 @@ Once we get into LinalgToNNVM passes we first do bufferize to generate Linalg on

Save the following mlir in /tmp/add.mlir
```mlir
func @add(%lhs: tensor<4xf32>, %rhs: tensor<4xf32>) -> tensor<4xf32>
attributes { iree.module.export } {
func @add(%lhs: tensor<4xf32>, %rhs: tensor<4xf32>) -> tensor<4xf32> {
%0 = "mhlo.add"(%lhs, %rhs) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
return %0 : tensor<4xf32>
}
Expand Down
13 changes: 5 additions & 8 deletions docs/developers/design_docs/simple_ir_walkthrough.md
Original file line number Diff line number Diff line change
Expand Up @@ -92,8 +92,7 @@ bridge can be used to convert the ops from GraphDef to XLA HLO, while a
based in MLIR is currently being written.

```mlir
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32>
attributes { iree.module.export } {
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> {
%0 = mhlo.multiply(%arg0, %arg1) : tensor<4xf32>
return %0 : tensor<4xf32>
}
Expand Down Expand Up @@ -127,8 +126,7 @@ semantics it's easy to use SSA use-def chains to ensure we are preserving the
expected behavior of the program.

```mlir
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32>
attributes {iree.module.export} {
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> {
%cst = constant dense<[4, 1, 1]> : tensor<3xi32>
%0 = iree.dispatch_region[%cst : tensor<3xi32>](%arg2 = %arg0 : tensor<4xf32>, %arg3 = %arg1 : tensor<4xf32>) : tensor<4xf32> {
%1 = mulf %arg2, %arg3 : tensor<4xf32>
Expand Down Expand Up @@ -178,8 +176,7 @@ module {
}
}
}
func @simple_mul(%arg0: memref<4xf32>, %arg1: memref<4xf32>) -> memref<4xf32>
attributes {iree.module.export} {
func @simple_mul(%arg0: memref<4xf32>, %arg1: memref<4xf32>) -> memref<4xf32> {
%0 = iree_interp.constant dense<[4, 1, 1]> : tensor<3xi32>
%1 = "iree_hl_seq.alloc_heap"() : () -> memref<4xf32>
iree_hl_seq.dispatch simple_mul_ex_dispatch_0::simple_mul_rgn_dispatch_0[%0 : memref<3xi32>](%arg0, %arg1, %1) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> ()
Expand Down Expand Up @@ -226,7 +223,7 @@ module {
}
}
func @simple_mul(%arg0: memref<4xf32>, %arg1: memref<4xf32>) -> memref<4xf32>
attributes {iree.module.export, iree.ordinal = 0 : i32} {
attributes {iree.ordinal = 0 : i32} {
%0 = "iree_ll_seq.alloc_heap"() : () -> memref<4xf32>
iree_ll_seq.static_dispatch simple_mul_ex_dispatch_0::simple_mul_rgn_dispatch_0[dense<[4, 1, 1]> : tensor<3xi32>](%arg0, %arg1, %0) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> ()
iree_ll_seq.return %0 : memref<4xf32>
Expand Down Expand Up @@ -333,7 +330,7 @@ module {
}
}
func @simple_mul(%arg0: memref<4xf32>, %arg1: memref<4xf32>) -> memref<4xf32>
attributes {iree.module.export, iree.ordinal = 0 : i32} {
attributes {iree.ordinal = 0 : i32} {
%0 = "iree_ll_seq.alloc_heap"() : () -> memref<4xf32>
iree_ll_seq.static_dispatch simple_mul_ex_dispatch_0::simple_mul_rgn_dispatch_0[dense<[4, 1, 1]> : tensor<3xi32>](%arg0, %arg1, %0) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> ()
iree_ll_seq.return %0 : memref<4xf32>
Expand Down
2 changes: 1 addition & 1 deletion docs/developers/design_roadmap.md
Original file line number Diff line number Diff line change
Expand Up @@ -135,7 +135,7 @@ comparison and increment on the device and insert a host readback to see if the
loop should continue:

```mlir
func @main() -> tensor<i32> attributes {iree.module.export, iree.reflection = {f = "I1!R6!B3!t6", fv = "1"}} {
func @main() -> tensor<i32> attributes {iree.reflection = {f = "I1!R6!B3!t6", fv = "1"}} {
%cst = constant dense<1> : tensor<i32>
%cst_0 = constant dense<3> : tensor<i32>
%cst_1 = constant dense<1> : vector<3xi32>
Expand Down
3 changes: 1 addition & 2 deletions docs/developers/developing_iree/profiling_vulkan_gpu.md
Original file line number Diff line number Diff line change
Expand Up @@ -27,8 +27,7 @@ For example, to package a module compiled from the following `mhlo-dot.mlir` as
an Android app:

```mlir
func @dot(%lhs: tensor<2x4xf32>, %rhs: tensor<4x2xf32>) -> tensor<2x2xf32>
attributes { iree.module.export } {
func @dot(%lhs: tensor<2x4xf32>, %rhs: tensor<4x2xf32>) -> tensor<2x2xf32> {
%0 = "mhlo.dot"(%lhs, %rhs) : (tensor<2x4xf32>, tensor<4x2xf32>) -> tensor<2x2xf32>
return %0 : tensor<2x2xf32>
}
Expand Down
11 changes: 5 additions & 6 deletions docs/developers/developing_iree/testing_guide.md
Original file line number Diff line number Diff line change
Expand Up @@ -222,37 +222,36 @@ returns no results and corresponds to a single test case.
As an example, here are some tests for the XLA HLO floor operation:

```mlir
func @tensor() attributes { iree.module.export } {
func @tensor() {
%input = iree.unfoldable_constant dense<[0.0, 1.1, 2.5, 4.9]> : tensor<4xf32>
%result = "mhlo.floor"(%input) : (tensor<4xf32>) -> tensor<4xf32>
check.expect_almost_eq_const(%result, dense<[0.0, 1.0, 2.0, 4.0]> : tensor<4xf32>): tensor<4xf32>
return
}
func @scalar() attributes { iree.module.export } {
func @scalar() {
%input = iree.unfoldable_constant dense<101.3> : tensor<f32>
%result = "mhlo.floor"(%input) : (tensor<f32>) -> tensor<f32>
check.expect_almost_eq_const(%result, dense<101.0> : tensor<f32>): tensor<f32>
return
}
func @double() attributes { iree.module.export } {
func @double() {
%input = iree.unfoldable_constant dense<11.2> : tensor<f64>
%result = "mhlo.floor"(%input) : (tensor<f64>) -> tensor<f64>
check.expect_almost_eq_const(%result, dense<11.0> : tensor<f64>): tensor<f64>
return
}
func @negative() attributes { iree.module.export } {
func @negative() {
%input = iree.unfoldable_constant dense<-1.1> : tensor<f32>
%result = "mhlo.floor"(%input) : (tensor<f32>) -> tensor<f32>
check.expect_almost_eq_const(%result, dense<-2.0> : tensor<f32>): tensor<f32>
return
}
```

The test case functions are exported using the `iree.module.export` attribute.
Each of these exported functions will be used to create a test case in gtest.
Test cases are created in gtest for each public function exported by the module.

Note the use of `iree.unfoldable_constant` to specify test constants. If we were
to use a regular constant, the compiler would "helpfully" fold away everything
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@ module @tf_module attributes {tf.versions = {bad_consumers = [], min_consumer =
// CHECK-NOT: tf._user_specified_name
// CHECK-NOT: tf._user_specified_name
// CHECK-NOT: tf._input_shapes
func @multiply__2_2__i32__uniform(%arg0: tensor<2xi32> {tf._user_specified_name = "args_0"}, %arg1: tensor<2xi32> {tf._user_specified_name = "args_1"}) -> tensor<2xi32> attributes {iree.module.export, iree.reflection = {abi = "sip", abiv = 1 : i32, sip = "I12!S9!k0_0k1_1R3!_0"}, tf._input_shapes = [#tf.shape<2>, #tf.shape<2>]} {
func @multiply__2_2__i32__uniform(%arg0: tensor<2xi32> {tf._user_specified_name = "args_0"}, %arg1: tensor<2xi32> {tf._user_specified_name = "args_1"}) -> tensor<2xi32> attributes {iree.reflection = {abi = "sip", abiv = 1 : i32, sip = "I12!S9!k0_0k1_1R3!_0"}, tf._input_shapes = [#tf.shape<2>, #tf.shape<2>]} {
// CHECK-NEXT: mhlo.multiply
%0 = mhlo.multiply %arg0, %arg1 : tensor<2xi32>
return %0 : tensor<2xi32>
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -53,6 +53,8 @@ class ConvertFunctionMetadataPass
// names to match up with those for readability so we extract them here.
// Is this ugly? Yeah - but such is what we have to deal with here.
void setupEntryPointAttrs(FuncOp funcOp, DictionaryAttr entryFunctionAttr) {
funcOp.setPublic();

auto inputsAttr =
entryFunctionAttr.get("inputs").template dyn_cast_or_null<StringAttr>();
auto outputsAttr = entryFunctionAttr.get("outputs")
Expand All @@ -64,8 +66,6 @@ class ConvertFunctionMetadataPass
return;
}

funcOp->setAttr("iree.module.export", UnitAttr::get(&getContext()));

SmallVector<std::string, 4> inputNames;
SmallVector<std::string, 4> outputNames;
splitFunctionIONames(inputsAttr, inputNames);
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -7,8 +7,6 @@ module attributes {tfl.schema_version = 3 : i32} {
// CHECK-SAME: ) -> (
// CHECK-SAME: tensor<?xf32> {iree.identifier = "output0"},
// CHECK-SAME: tensor<?xf32> {iree.identifier = "output1"})
// CHECK-SAME: attributes
// CHECK-SAME: iree.module.export
func @main(%arg0: tensor<?xf32>, %arg1: tensor<?xf32>) -> (tensor<?xf32>, tensor<?xf32>) attributes {
tf.entry_function = {inputs = "input0,input1", outputs = "output0,output1"}
} {
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -229,13 +229,13 @@ int main(int argc, char **argv) {
// function.
std::string entryName = "main";
SymbolTable symbolTable(module.get());
Operation *mainFunc = symbolTable.lookup(entryName);
auto mainFunc = symbolTable.lookup<FuncOp>(entryName);
if (!mainFunc) {
llvm::errs() << "Unable to find main function '" << entryName
<< "' in converted module.\n";
return 3;
}
mainFunc->setAttr("iree.module.export", UnitAttr::get(&context));
mainFunc.setPublic();

// Save.
auto saveToFile = [&](llvm::StringRef savePath) -> LogicalResult {
Expand Down
21 changes: 0 additions & 21 deletions iree/base/BUILD
Original file line number Diff line number Diff line change
Expand Up @@ -76,27 +76,6 @@ cc_library(
],
)

cc_library(
name = "signature_parser",
srcs = ["signature_parser.cc"],
hdrs = ["signature_parser.h"],
deps = [
"@com_google_absl//absl/strings",
"@com_google_absl//absl/types:optional",
"@com_google_absl//absl/types:span",
],
)

cc_test(
name = "signature_parser_test",
srcs = ["signature_parser_test.cc"],
deps = [
":signature_parser",
"//iree/testing:gtest",
"//iree/testing:gtest_main",
],
)

cc_library(
name = "status",
hdrs = ["status.h"],
Expand Down
25 changes: 0 additions & 25 deletions iree/base/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -59,31 +59,6 @@ iree_cc_library(
PUBLIC
)

iree_cc_library(
NAME
signature_parser
HDRS
"signature_parser.h"
SRCS
"signature_parser.cc"
DEPS
absl::optional
absl::span
absl::strings
PUBLIC
)

iree_cc_test(
NAME
signature_parser_test
SRCS
"signature_parser_test.cc"
DEPS
::signature_parser
iree::testing::gtest
iree::testing::gtest_main
)

iree_cc_library(
NAME
status
Expand Down
Loading

0 comments on commit 47be229

Please sign in to comment.