Created
March 30, 2020 03:53
-
-
Save stellaraccident/58547467dd855acd3e09e78c4f336cc6 to your computer and use it in GitHub Desktop.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| *** IR Dump Before Canonicalizer *** | |
| module { | |
| func @simple_mul(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> attributes {iree.module.export} { | |
| %0 = "xla_hlo.abs"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| return %0 : tensor<?x?xf32> | |
| } | |
| } | |
| *** IR Dump Before xla_hlo::(anonymous namespace)::LegalizeControlFlow *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> attributes {iree.module.export} { | |
| %0 = "xla_hlo.abs"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| return %0 : tensor<?x?xf32> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::(anonymous namespace)::FlattenTuplesInCFGPass *** | |
| module { | |
| func @simple_mul(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> attributes {iree.module.export} { | |
| %0 = "xla_hlo.abs"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| return %0 : tensor<?x?xf32> | |
| } | |
| } | |
| *** IR Dump Before InlinerPass *** | |
| module { | |
| func @simple_mul(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> attributes {iree.module.export} { | |
| %0 = "xla_hlo.abs"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| return %0 : tensor<?x?xf32> | |
| } | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> attributes {iree.module.export} { | |
| %0 = "xla_hlo.abs"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| return %0 : tensor<?x?xf32> | |
| } | |
| *** IR Dump Before CSE *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> attributes {iree.module.export} { | |
| %0 = "xla_hlo.abs"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| return %0 : tensor<?x?xf32> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::LegalizeInputTypesPass *** | |
| module { | |
| func @simple_mul(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> attributes {iree.module.export} { | |
| %0 = "xla_hlo.abs"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| return %0 : tensor<?x?xf32> | |
| } | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::MaterializeExportedReflectionPass *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> attributes {iree.module.export} { | |
| %0 = "xla_hlo.abs"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| return %0 : tensor<?x?xf32> | |
| } | |
| *** IR Dump Before iree_compiler::Shape::(anonymous namespace)::ExpandFunctionDynamicDimsPass *** | |
| func @simple_mul(%arg0: tensor<?x?xf32> {iree.reflection = {f_partial = "I10!B7!d-1d-1"}}) -> (tensor<?x?xf32> {iree.reflection = {f_partial = "R10!B7!d-1d-1"}}) attributes {iree.module.export} { | |
| %0 = "xla_hlo.abs"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| return %0 : tensor<?x?xf32> | |
| } | |
| *** IR Dump Before iree_compiler::Shape::(anonymous namespace)::TieDynamicShapesPass *** | |
| func @simple_mul(%arg0: tensor<?x?xf32> {iree.reflection = {f_partial = "I10!B7!d-1d-1"}}, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32> {iree.reflection = {f_partial = "R10!B7!d-1d-1"}}, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export} { | |
| %0 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %1 = "xla_hlo.abs"(%0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %2 = shapex.get_ranked_shape %1 : tensor<?x?xf32> -> !shapex.ranked_shape<[?,?]> | |
| return %1, %2 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before iree_compiler::Shape::(anonymous namespace)::MaterializeShapeCalculationsPass *** | |
| func @simple_mul(%arg0: tensor<?x?xf32> {iree.reflection = {f_partial = "I10!B7!d-1d-1"}}, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32> {iree.reflection = {f_partial = "R10!B7!d-1d-1"}}, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export} { | |
| %0 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %1 = "xla_hlo.abs"(%0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %2 = shapex.get_ranked_shape %1 : tensor<?x?xf32> -> !shapex.ranked_shape<[?,?]> | |
| %3 = shapex.tie_shape %1, %2 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %4 = shapex.get_ranked_shape %3 : tensor<?x?xf32> -> !shapex.ranked_shape<[?,?]> | |
| return %3, %4 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::MergeExportedReflectionPass *** | |
| func @simple_mul(%arg0: tensor<?x?xf32> {iree.reflection = {f_partial = "I10!B7!d-1d-1"}}, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32> {iree.reflection = {f_partial = "R10!B7!d-1d-1"}}, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export} { | |
| %0 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %1 = "xla_hlo.abs"(%0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %2 = shapex.tie_shape %1, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %2, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::HLOToHLOPreprocessing *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %1 = "xla_hlo.abs"(%0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %2 = shapex.tie_shape %1, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %2, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::PrePartitioningConversionPass *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %1 = "xla_hlo.abs"(%0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %2 = shapex.tie_shape %1, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %2, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::DispatchabilityAnalysisPass *** | |
| module { | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %1 = "xla_hlo.abs"(%0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %2 = shapex.tie_shape %1, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %2, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::IdentifyDispatchRegionsPass *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %1 = "xla_hlo.abs"(%0) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %2 = shapex.tie_shape %1, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %2, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before CSE *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %4 = flow.dispatch.region[%2 : index](%arg2 = %3 : tensor<?x?xf32>, %arg3 = %arg1 : !shapex.ranked_shape<[?,?]>) -> tensor<?x?xf32> { | |
| %6 = shapex.tie_shape %arg2, %arg3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %7 = "xla_hlo.abs"(%6) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %8 = shapex.tie_shape %7, %arg3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| flow.return %8 : tensor<?x?xf32> | |
| } | |
| %5 = shapex.tie_shape %4, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %5, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::FoldCompatibleDispatchRegionsPass *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %4 = flow.dispatch.region[%2 : index](%arg2 = %3 : tensor<?x?xf32>, %arg3 = %arg1 : !shapex.ranked_shape<[?,?]>) -> tensor<?x?xf32> { | |
| %6 = shapex.tie_shape %arg2, %arg3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %7 = "xla_hlo.abs"(%6) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %8 = shapex.tie_shape %7, %arg3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| flow.return %8 : tensor<?x?xf32> | |
| } | |
| %5 = shapex.tie_shape %4, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %5, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::RematerializeDispatchConstantsPass *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %4 = flow.dispatch.region[%2 : index](%arg2 = %3 : tensor<?x?xf32>, %arg3 = %arg1 : !shapex.ranked_shape<[?,?]>) -> tensor<?x?xf32> { | |
| %6 = shapex.tie_shape %arg2, %arg3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %7 = "xla_hlo.abs"(%6) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %8 = shapex.tie_shape %7, %arg3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| flow.return %8 : tensor<?x?xf32> | |
| } | |
| %5 = shapex.tie_shape %4, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %5, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::OutlineDispatchRegionsPass *** | |
| module { | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %4 = flow.dispatch.region[%2 : index](%arg2 = %3 : tensor<?x?xf32>, %arg3 = %arg1 : !shapex.ranked_shape<[?,?]>) -> tensor<?x?xf32> { | |
| %6 = shapex.tie_shape %arg2, %arg3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %7 = "xla_hlo.abs"(%6) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %8 = shapex.tie_shape %7, %arg3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| flow.return %8 : tensor<?x?xf32> | |
| } | |
| %5 = shapex.tie_shape %4, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %5, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %4 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %5 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %6 = flow.dispatch @simple_mul_ex_dispatch_0::@simple_mul_ex_dispatch_0[%2 : index](%3, %4, %5) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| %7 = shapex.tie_shape %6, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %7, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::PostPartitioningConversionPass *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %4 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %5 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %6 = flow.dispatch @simple_mul_ex_dispatch_0::@simple_mul_ex_dispatch_0[%2 : index](%3, %4, %5) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| %7 = shapex.tie_shape %6, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %7, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::Flow::FormStreamsPass *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %4 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %5 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %6 = flow.dispatch @simple_mul_ex_dispatch_0::@simple_mul_ex_dispatch_0[%2 : index](%3, %4, %5) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| %7 = shapex.tie_shape %6, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %7, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before CSE *** | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %4 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %5 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %6 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %7 = flow.ex.stream.fragment(%arg2 = %arg0 : tensor<?x?xf32>, %arg3 = %5 : index, %arg4 = %6 : index, %arg5 = %2 : index, %arg6 = %3 : index, %arg7 = %4 : index) -> tensor<?x?xf32> { | |
| %8 = shapex.make_ranked_shape %arg3, %arg4 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %9 = shapex.tie_shape %arg2, %8 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %10 = flow.dispatch @simple_mul_ex_dispatch_0::@simple_mul_ex_dispatch_0[%arg5 : index](%9, %arg6, %arg7) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| %11 = shapex.tie_shape %10, %8 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| flow.return %11 : tensor<?x?xf32> | |
| } | |
| return %7, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| module { | |
| flow.executable @simple_mul_ex_dispatch_0 { | |
| flow.dispatch.entry @simple_mul_ex_dispatch_0 | |
| module { | |
| func @simple_mul_ex_dispatch_0(%arg0: tensor<?x?xf32>, %arg1: index, %arg2: index) -> tensor<?x?xf32> { | |
| %0 = shapex.make_ranked_shape %arg1, %arg2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %1 = shapex.tie_shape %arg0, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %2 = "xla_hlo.abs"(%1) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %3 = shapex.tie_shape %2, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %3 : tensor<?x?xf32> | |
| } | |
| } | |
| } | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = flow.ex.stream.fragment(%arg2 = %arg0 : tensor<?x?xf32>, %arg3 = %0 : index, %arg4 = %1 : index, %arg5 = %2 : index, %arg6 = %0 : index, %arg7 = %1 : index) -> tensor<?x?xf32> { | |
| %4 = shapex.make_ranked_shape %arg3, %arg4 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %5 = shapex.tie_shape %arg2, %4 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %6 = flow.dispatch @simple_mul_ex_dispatch_0::@simple_mul_ex_dispatch_0[%arg5 : index](%5, %arg6, %arg7) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| %7 = shapex.tie_shape %6, %4 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| flow.return %7 : tensor<?x?xf32> | |
| } | |
| return %3, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| } | |
| *** IR Dump Before iree_compiler::IREE::HAL::MaterializeInterfacesPass *** | |
| module { | |
| flow.executable @simple_mul_ex_dispatch_0 { | |
| flow.dispatch.entry @simple_mul_ex_dispatch_0 | |
| module { | |
| func @simple_mul_ex_dispatch_0(%arg0: tensor<?x?xf32>, %arg1: index, %arg2: index) -> tensor<?x?xf32> { | |
| %0 = shapex.make_ranked_shape %arg1, %arg2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %1 = shapex.tie_shape %arg0, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %2 = "xla_hlo.abs"(%1) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %3 = shapex.tie_shape %2, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %3 : tensor<?x?xf32> | |
| } | |
| } | |
| } | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = flow.ex.stream.fragment(%arg2 = %arg0 : tensor<?x?xf32>, %arg3 = %0 : index, %arg4 = %1 : index, %arg5 = %2 : index) -> tensor<?x?xf32> { | |
| %4 = shapex.make_ranked_shape %arg3, %arg4 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %5 = shapex.tie_shape %arg2, %4 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %6 = flow.dispatch @simple_mul_ex_dispatch_0::@simple_mul_ex_dispatch_0[%arg5 : index](%5, %arg3, %arg4) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| %7 = shapex.tie_shape %6, %4 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| flow.return %7 : tensor<?x?xf32> | |
| } | |
| return %3, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| } | |
| *** IR Dump Before iree_compiler::IREE::HAL::TranslateExecutablesPass *** | |
| module { | |
| hal.executable @simple_mul_ex_dispatch_0 attributes {sym_visibility = "private"} { | |
| hal.interface @legacy_io attributes {push_constants = 2 : i32} { | |
| hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read" | |
| hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" | |
| } | |
| hal.executable.entry_point @simple_mul_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<?x?xf32>, index, index) -> tensor<?x?xf32>, workgroup_size = [1 : index, 1 : index, 1 : index]} | |
| hal.executable.source { | |
| module { | |
| flow.executable @simple_mul_ex_dispatch_0 { | |
| flow.dispatch.entry @simple_mul_ex_dispatch_0 | |
| module { | |
| func @simple_mul_ex_dispatch_0() { | |
| %c0 = constant 0 : index | |
| %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<?x?xf32> | |
| %1 = hal.interface.load.constant offset = 0 : index | |
| %2 = hal.interface.load.constant offset = 1 : index | |
| %3 = call @simple_mul_ex_dispatch_0_impl(%0, %1, %2) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 : tensor<?x?xf32> | |
| return | |
| } | |
| func @simple_mul_ex_dispatch_0_impl(%arg0: tensor<?x?xf32>, %arg1: index, %arg2: index) -> tensor<?x?xf32> attributes {sym_visibility = "private"} { | |
| %0 = shapex.make_ranked_shape %arg1, %arg2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %1 = shapex.tie_shape %arg0, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %2 = "xla_hlo.abs"(%1) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %3 = shapex.tie_shape %2, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %3 : tensor<?x?xf32> | |
| } | |
| hal.interface @legacy_io attributes {push_constants = 2 : i32, sym_visibility = "private"} { | |
| hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read" | |
| hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" | |
| } | |
| } | |
| } | |
| } | |
| } | |
| } | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = flow.ex.stream.fragment(%arg2 = %arg0 : tensor<?x?xf32>, %arg3 = %0 : index, %arg4 = %1 : index, %arg5 = %2 : index) -> tensor<?x?xf32> { | |
| %4 = shapex.make_ranked_shape %arg3, %arg4 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %5 = shapex.tie_shape %arg2, %4 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %6 = flow.dispatch @simple_mul_ex_dispatch_0::@simple_mul_ex_dispatch_0[%arg5 : index](%5, %arg3, %arg4) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| %7 = shapex.tie_shape %6, %4 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| flow.return %7 : tensor<?x?xf32> | |
| } | |
| return %3, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| module { | |
| func @simple_mul_ex_dispatch_0() attributes {iree.module.export} { | |
| %c0 = constant 0 : index | |
| %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<?x?xf32> | |
| %1 = hal.interface.load.constant offset = 0 : index | |
| %2 = hal.interface.load.constant offset = 1 : index | |
| %3 = call @simple_mul_ex_dispatch_0_impl(%0, %1, %2) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 : tensor<?x?xf32> | |
| return | |
| } | |
| func @simple_mul_ex_dispatch_0_impl(%arg0: tensor<?x?xf32>, %arg1: index, %arg2: index) -> tensor<?x?xf32> attributes {sym_visibility = "private"} { | |
| %0 = shapex.make_ranked_shape %arg1, %arg2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %1 = shapex.tie_shape %arg0, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %2 = "xla_hlo.abs"(%1) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %3 = shapex.tie_shape %2, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %3 : tensor<?x?xf32> | |
| } | |
| hal.interface @legacy_io attributes {push_constants = 2 : i32, sym_visibility = "private"} { | |
| hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read" | |
| hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" | |
| } | |
| } | |
| *** IR Dump Before xla_hlo::(anonymous namespace)::LegalizeControlFlow *** | |
| func @simple_mul_ex_dispatch_0() attributes {iree.module.export} { | |
| %c0 = constant 0 : index | |
| %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<?x?xf32> | |
| %1 = hal.interface.load.constant offset = 0 : index | |
| %2 = hal.interface.load.constant offset = 1 : index | |
| %3 = call @simple_mul_ex_dispatch_0_impl(%0, %1, %2) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 : tensor<?x?xf32> | |
| return | |
| } | |
| *** IR Dump Before xla_hlo::(anonymous namespace)::LegalizeControlFlow *** | |
| func @simple_mul_ex_dispatch_0_impl(%arg0: tensor<?x?xf32>, %arg1: index, %arg2: index) -> tensor<?x?xf32> attributes {sym_visibility = "private"} { | |
| %0 = shapex.make_ranked_shape %arg1, %arg2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %1 = shapex.tie_shape %arg0, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %2 = "xla_hlo.abs"(%1) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %3 = shapex.tie_shape %2, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %3 : tensor<?x?xf32> | |
| } | |
| *** IR Dump Before InlinerPass *** | |
| module { | |
| func @simple_mul_ex_dispatch_0() attributes {iree.module.export} { | |
| %c0 = constant 0 : index | |
| %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<?x?xf32> | |
| %1 = hal.interface.load.constant offset = 0 : index | |
| %2 = hal.interface.load.constant offset = 1 : index | |
| %3 = call @simple_mul_ex_dispatch_0_impl(%0, %1, %2) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 : tensor<?x?xf32> | |
| return | |
| } | |
| func @simple_mul_ex_dispatch_0_impl(%arg0: tensor<?x?xf32>, %arg1: index, %arg2: index) -> tensor<?x?xf32> attributes {sym_visibility = "private"} { | |
| %0 = shapex.make_ranked_shape %arg1, %arg2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %1 = shapex.tie_shape %arg0, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %2 = "xla_hlo.abs"(%1) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %3 = shapex.tie_shape %2, %0 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| return %3 : tensor<?x?xf32> | |
| } | |
| hal.interface @legacy_io attributes {push_constants = 2 : i32, sym_visibility = "private"} { | |
| hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read" | |
| hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" | |
| } | |
| } | |
| *** IR Dump Before SymbolDCE *** | |
| module { | |
| func @simple_mul_ex_dispatch_0() attributes {iree.module.export} { | |
| %c0 = constant 0 : index | |
| %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<?x?xf32> | |
| %1 = hal.interface.load.constant offset = 0 : index | |
| %2 = hal.interface.load.constant offset = 1 : index | |
| %3 = shapex.make_ranked_shape %1, %2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %4 = shapex.tie_shape %0, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %5 = "xla_hlo.abs"(%4) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %6 = shapex.tie_shape %5, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| hal.interface.store.tensor %6, @legacy_io::@ret0, offset = %c0 : tensor<?x?xf32> | |
| return | |
| } | |
| hal.interface @legacy_io attributes {push_constants = 2 : i32, sym_visibility = "private"} { | |
| hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read" | |
| hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" | |
| } | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| func @simple_mul_ex_dispatch_0() attributes {iree.module.export} { | |
| %c0 = constant 0 : index | |
| %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<?x?xf32> | |
| %1 = hal.interface.load.constant offset = 0 : index | |
| %2 = hal.interface.load.constant offset = 1 : index | |
| %3 = shapex.make_ranked_shape %1, %2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %4 = shapex.tie_shape %0, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %5 = "xla_hlo.abs"(%4) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %6 = shapex.tie_shape %5, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| hal.interface.store.tensor %6, @legacy_io::@ret0, offset = %c0 : tensor<?x?xf32> | |
| return | |
| } | |
| *** IR Dump Before CSE *** | |
| func @simple_mul_ex_dispatch_0() attributes {iree.module.export} { | |
| %c0 = constant 0 : index | |
| %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<?x?xf32> | |
| %1 = hal.interface.load.constant offset = 0 : index | |
| %2 = hal.interface.load.constant offset = 1 : index | |
| %3 = shapex.make_ranked_shape %1, %2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %4 = shapex.tie_shape %0, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %5 = "xla_hlo.abs"(%4) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %6 = shapex.tie_shape %5, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| hal.interface.store.tensor %6, @legacy_io::@ret0, offset = %c0 : tensor<?x?xf32> | |
| return | |
| } | |
| *** IR Dump Before iree_compiler::IREE::VMLA::UnrollReductionsPass *** | |
| func @simple_mul_ex_dispatch_0() attributes {iree.module.export} { | |
| %c0 = constant 0 : index | |
| %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<?x?xf32> | |
| %1 = hal.interface.load.constant offset = 0 : index | |
| %2 = hal.interface.load.constant offset = 1 : index | |
| %3 = shapex.make_ranked_shape %1, %2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %4 = shapex.tie_shape %0, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %5 = "xla_hlo.abs"(%4) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %6 = shapex.tie_shape %5, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| hal.interface.store.tensor %6, @legacy_io::@ret0, offset = %c0 : tensor<?x?xf32> | |
| return | |
| } | |
| *** IR Dump Before CSE *** | |
| func @simple_mul_ex_dispatch_0() attributes {iree.module.export} { | |
| %c0 = constant 0 : index | |
| %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<?x?xf32> | |
| %1 = hal.interface.load.constant offset = 0 : index | |
| %2 = hal.interface.load.constant offset = 1 : index | |
| %3 = shapex.make_ranked_shape %1, %2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %4 = shapex.tie_shape %0, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %5 = "xla_hlo.abs"(%4) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %6 = shapex.tie_shape %5, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| hal.interface.store.tensor %6, @legacy_io::@ret0, offset = %c0 : tensor<?x?xf32> | |
| return | |
| } | |
| *** IR Dump Before iree_compiler::IREE::VMLA::ConversionPass *** | |
| module { | |
| func @simple_mul_ex_dispatch_0() attributes {iree.module.export} { | |
| %c0 = constant 0 : index | |
| %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<?x?xf32> | |
| %1 = hal.interface.load.constant offset = 0 : index | |
| %2 = hal.interface.load.constant offset = 1 : index | |
| %3 = shapex.make_ranked_shape %1, %2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %4 = shapex.tie_shape %0, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %5 = "xla_hlo.abs"(%4) : (tensor<?x?xf32>) -> tensor<?x?xf32> | |
| %6 = shapex.tie_shape %5, %3 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| hal.interface.store.tensor %6, @legacy_io::@ret0, offset = %c0 : tensor<?x?xf32> | |
| return | |
| } | |
| hal.interface @legacy_io attributes {push_constants = 2 : i32, sym_visibility = "private"} { | |
| hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read" | |
| hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" | |
| } | |
| } | |
| *** IR Dump Before CSE *** | |
| func @simple_mul_ex_dispatch_0(%arg0: !vmla.interface) attributes {iree.module.export} { | |
| %c4 = constant 4 : index | |
| %c0 = constant 0 : index | |
| %c4_0 = constant 4 : index | |
| %0 = "vmla.interface.const"(%arg0) {offset = 1 : index} : (!vmla.interface) -> index | |
| %1 = "vmla.interface.const"(%arg0) {offset = 0 : index} : (!vmla.interface) -> index | |
| %c4_1 = constant 4 : index | |
| %c0_2 = constant 0 : index | |
| %2 = "vmla.interface.binding"(%arg0) {binding = 0 : i32, set = 0 : i32} : (!vmla.interface) -> !vmla.buffer | |
| %3 = shapex.make_ranked_shape %1, %0 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %4 = muli %c4, %1 : index | |
| %5 = muli %4, %0 : index | |
| %6 = muli %c4_0, %1 : index | |
| %7 = muli %6, %0 : index | |
| %8 = "vmla.buffer.alloc"(%7) : (index) -> !vmla.buffer | |
| %9 = muli %c4_1, %1 : index | |
| %10 = muli %9, %0 : index | |
| %11 = "vmla.buffer.view"(%2, %c0_2, %10) : (!vmla.buffer, index, index) -> !vmla.buffer | |
| %12 = shapex.tie_shape %11, %3 : !vmla.buffer, !shapex.ranked_shape<[?,?]> | |
| "vmla.abs"(%12, %8) {element_type = f32} : (!vmla.buffer, !vmla.buffer) -> () | |
| %13 = shapex.tie_shape %8, %3 : !vmla.buffer, !shapex.ranked_shape<[?,?]> | |
| %14 = "vmla.interface.binding"(%arg0) {binding = 1 : i32, set = 0 : i32} : (!vmla.interface) -> !vmla.buffer | |
| "vmla.buffer.copy"(%13, %c0, %14, %c0_2, %5) : (!vmla.buffer, index, !vmla.buffer, index, index) -> () | |
| return | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| func @simple_mul_ex_dispatch_0(%arg0: !vmla.interface) attributes {iree.module.export} { | |
| %c4 = constant 4 : index | |
| %c0 = constant 0 : index | |
| %0 = "vmla.interface.const"(%arg0) {offset = 1 : index} : (!vmla.interface) -> index | |
| %1 = "vmla.interface.const"(%arg0) {offset = 0 : index} : (!vmla.interface) -> index | |
| %2 = "vmla.interface.binding"(%arg0) {binding = 0 : i32, set = 0 : i32} : (!vmla.interface) -> !vmla.buffer | |
| %3 = shapex.make_ranked_shape %1, %0 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %4 = muli %c4, %1 : index | |
| %5 = muli %4, %0 : index | |
| %6 = "vmla.buffer.alloc"(%5) : (index) -> !vmla.buffer | |
| %7 = "vmla.buffer.view"(%2, %c0, %5) : (!vmla.buffer, index, index) -> !vmla.buffer | |
| %8 = shapex.tie_shape %7, %3 : !vmla.buffer, !shapex.ranked_shape<[?,?]> | |
| "vmla.abs"(%8, %6) {element_type = f32} : (!vmla.buffer, !vmla.buffer) -> () | |
| %9 = shapex.tie_shape %6, %3 : !vmla.buffer, !shapex.ranked_shape<[?,?]> | |
| %10 = "vmla.interface.binding"(%arg0) {binding = 1 : i32, set = 0 : i32} : (!vmla.interface) -> !vmla.buffer | |
| "vmla.buffer.copy"(%9, %c0, %10, %c0, %5) : (!vmla.buffer, index, !vmla.buffer, index, index) -> () | |
| return | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| module { | |
| func @simple_mul_ex_dispatch_0(%arg0: !vmla.interface) attributes {iree.module.export} { | |
| %c4 = constant 4 : index | |
| %c0 = constant 0 : index | |
| %0 = "vmla.interface.const"(%arg0) {offset = 1 : index} : (!vmla.interface) -> index | |
| %1 = "vmla.interface.const"(%arg0) {offset = 0 : index} : (!vmla.interface) -> index | |
| %2 = "vmla.interface.binding"(%arg0) {binding = 0 : i32, set = 0 : i32} : (!vmla.interface) -> !vmla.buffer | |
| %3 = shapex.make_ranked_shape %1, %0 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %4 = muli %1, %c4 : index | |
| %5 = muli %4, %0 : index | |
| %6 = "vmla.buffer.alloc"(%5) : (index) -> !vmla.buffer | |
| %7 = "vmla.buffer.view"(%2, %c0, %5) : (!vmla.buffer, index, index) -> !vmla.buffer | |
| %8 = shapex.tie_shape %7, %3 : !vmla.buffer, !shapex.ranked_shape<[?,?]> | |
| "vmla.abs"(%8, %6) {element_type = f32} : (!vmla.buffer, !vmla.buffer) -> () | |
| %9 = shapex.tie_shape %6, %3 : !vmla.buffer, !shapex.ranked_shape<[?,?]> | |
| %10 = "vmla.interface.binding"(%arg0) {binding = 1 : i32, set = 0 : i32} : (!vmla.interface) -> !vmla.buffer | |
| "vmla.buffer.copy"(%9, %c0, %10, %c0, %5) : (!vmla.buffer, index, !vmla.buffer, index, index) -> () | |
| return | |
| } | |
| } | |
| *** IR Dump Before iree_compiler::IREE::VM::ConversionPass *** | |
| module { | |
| func @simple_mul_ex_dispatch_0(%arg0: !vmla.interface) attributes {iree.module.export} { | |
| %c4 = constant 4 : index | |
| %c0 = constant 0 : index | |
| %0 = "vmla.interface.const"(%arg0) {offset = 1 : index} : (!vmla.interface) -> index | |
| %1 = "vmla.interface.const"(%arg0) {offset = 0 : index} : (!vmla.interface) -> index | |
| %2 = "vmla.interface.binding"(%arg0) {binding = 0 : i32, set = 0 : i32} : (!vmla.interface) -> !vmla.buffer | |
| %3 = shapex.make_ranked_shape %1, %0 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %4 = muli %1, %c4 : index | |
| %5 = muli %4, %0 : index | |
| %6 = "vmla.buffer.alloc"(%5) : (index) -> !vmla.buffer | |
| %7 = "vmla.buffer.view"(%2, %c0, %5) : (!vmla.buffer, index, index) -> !vmla.buffer | |
| %8 = shapex.tie_shape %7, %3 : !vmla.buffer, !shapex.ranked_shape<[?,?]> | |
| "vmla.abs"(%8, %6) {element_type = f32} : (!vmla.buffer, !vmla.buffer) -> () | |
| %9 = shapex.tie_shape %6, %3 : !vmla.buffer, !shapex.ranked_shape<[?,?]> | |
| %10 = "vmla.interface.binding"(%arg0) {binding = 1 : i32, set = 0 : i32} : (!vmla.interface) -> !vmla.buffer | |
| "vmla.buffer.copy"(%9, %c0, %10, %c0, %5) : (!vmla.buffer, index, !vmla.buffer, index, index) -> () | |
| return | |
| } | |
| } | |
| *** IR Dump Before iree_compiler::IREE::VM::GlobalInitializationPass *** | |
| vm.module @module { | |
| vm.func @simple_mul_ex_dispatch_0(%arg0: !vm.ref<!vmla.interface>) { | |
| %c4 = vm.const.i32 4 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %c1 = vm.const.i32 1 : i32 | |
| %0 = vm.call @vmla.interface.const(%arg0, %c1) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %zero_0 = vm.const.i32.zero : i32 | |
| %1 = vm.call @vmla.interface.const(%arg0, %zero_0) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %zero_1 = vm.const.i32.zero : i32 | |
| %zero_2 = vm.const.i32.zero : i32 | |
| %ref = vm.call @vmla.interface.binding(%arg0, %zero_1, %zero_2) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| %2 = vm.mul.i32 %1, %c4 : i32 | |
| %3 = vm.mul.i32 %2, %0 : i32 | |
| %ref_3 = vm.call @vmla.buffer.alloc(%3) : (i32) -> !vm.ref<!vmla.buffer> | |
| %ref_4 = vm.call @vmla.buffer.view(%ref, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.abs.f32(%ref_4, %ref_3) : (!vm.ref<!vmla.buffer>, !vm.ref<!vmla.buffer>) -> () | |
| %zero_5 = vm.const.i32.zero : i32 | |
| %c1_6 = vm.const.i32 1 : i32 | |
| %ref_7 = vm.call @vmla.interface.binding(%arg0, %zero_5, %c1_6) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.buffer.copy(%ref_3, %zero, %ref_7, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, !vm.ref<!vmla.buffer>, i32, i32) -> () | |
| vm.return | |
| } | |
| vm.export @simple_mul_ex_dispatch_0 | |
| vm.import @vmla.interface.current() -> !vm.ref<!vmla.interface> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.interface.const(%interface : !vm.ref<!vmla.interface>, %offset : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.interface.binding(%interface : !vm.ref<!vmla.interface>, %set : i32, %binding : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.const(%value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.alloc(%byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.clone(%src : !vm.ref<!vmla.buffer>) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.byte_length(%value : !vm.ref<!vmla.buffer>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.view(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32, %byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.copy(%src : !vm.ref<!vmla.buffer>, %src_byte_offset : i32, %dst : !vm.ref<!vmla.buffer>, %dst_byte_offset : i32, %byte_length : i32) attributes {sym_visibility = "private"} | |
| vm.import @vmla.buffer.fill(%value : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.buffer.load.i32(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.cmp.i8(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cmp.i16(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cmp.i32(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cmp.f32(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.select.x8(%cond : !vm.ref<!vmla.buffer>, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.select.x16(%cond : !vm.ref<!vmla.buffer>, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.select.x32(%cond : !vm.ref<!vmla.buffer>, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.copy.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %src_indices : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %dst_indices : i32..., %lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.copy.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %src_indices : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %dst_indices : i32..., %lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.copy.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %src_indices : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %dst_indices : i32..., %lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.transpose.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %permutation : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.transpose.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %permutation : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.transpose.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %permutation : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reverse.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dimensions : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reverse.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dimensions : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reverse.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dimensions : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pad.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %value : !vm.ref<!vmla.buffer>, %value_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %edge_padding_low : i32..., %edge_padding_high : i32..., %interior_padding : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pad.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %value : !vm.ref<!vmla.buffer>, %value_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %edge_padding_low : i32..., %edge_padding_high : i32..., %interior_padding : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pad.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %value : !vm.ref<!vmla.buffer>, %value_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %edge_padding_low : i32..., %edge_padding_high : i32..., %interior_padding : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.broadcast.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.broadcast.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.broadcast.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tile.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tile.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tile.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.not.x8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.not.x16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.not.x32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.and.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.and.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.and.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.or.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.or.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.or.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.xor.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.xor.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.xor.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shl.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shl.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shl.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.u8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.u16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.u32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.u8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.u16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.u32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.u8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.u16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.u32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pow.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.exp.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.log.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rsqrt.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sqrt.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cos.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sin.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tanh.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.atan2.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.i8(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.i16(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.i32(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.f32(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.floor.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.ceil.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i8.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i8.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i8.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i16.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i16.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i16.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i32.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i32.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i32.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.f32.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.f32.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.f32.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.conv.f32f32.f32(%input : !vm.ref<!vmla.buffer>, %input_shape : i32..., %filter : !vm.ref<!vmla.buffer>, %filter_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %window_strides : i32..., %padding : i32..., %lhs_dilation : i32..., %rhs_dilation : i32..., %feature_group_count : i32, %batch_group_count : i32) attributes {sym_visibility = "private"} | |
| vm.import @vmla.batch.matmul.f32f32.f32(%lhs : !vm.ref<!vmla.buffer>, %lhs_shape : i32..., %rhs : !vm.ref<!vmla.buffer>, %rhs_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.i8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.i16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.i32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.f32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.i8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.i16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.i32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.f32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.i8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.i16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.i32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.f32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| } | |
| *** IR Dump Before CSE *** | |
| module { | |
| vm.module @module { | |
| vm.func @simple_mul_ex_dispatch_0(%arg0: !vm.ref<!vmla.interface>) { | |
| %c4 = vm.const.i32 4 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %c1 = vm.const.i32 1 : i32 | |
| %0 = vm.call @vmla.interface.const(%arg0, %c1) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %zero_0 = vm.const.i32.zero : i32 | |
| %1 = vm.call @vmla.interface.const(%arg0, %zero_0) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %zero_1 = vm.const.i32.zero : i32 | |
| %zero_2 = vm.const.i32.zero : i32 | |
| %ref = vm.call @vmla.interface.binding(%arg0, %zero_1, %zero_2) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| %2 = vm.mul.i32 %1, %c4 : i32 | |
| %3 = vm.mul.i32 %2, %0 : i32 | |
| %ref_3 = vm.call @vmla.buffer.alloc(%3) : (i32) -> !vm.ref<!vmla.buffer> | |
| %ref_4 = vm.call @vmla.buffer.view(%ref, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.abs.f32(%ref_4, %ref_3) : (!vm.ref<!vmla.buffer>, !vm.ref<!vmla.buffer>) -> () | |
| %zero_5 = vm.const.i32.zero : i32 | |
| %c1_6 = vm.const.i32 1 : i32 | |
| %ref_7 = vm.call @vmla.interface.binding(%arg0, %zero_5, %c1_6) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.buffer.copy(%ref_3, %zero, %ref_7, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, !vm.ref<!vmla.buffer>, i32, i32) -> () | |
| vm.return | |
| } | |
| vm.export @simple_mul_ex_dispatch_0 | |
| vm.import @vmla.interface.current() -> !vm.ref<!vmla.interface> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.interface.const(%interface : !vm.ref<!vmla.interface>, %offset : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.interface.binding(%interface : !vm.ref<!vmla.interface>, %set : i32, %binding : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.const(%value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.alloc(%byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.clone(%src : !vm.ref<!vmla.buffer>) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.byte_length(%value : !vm.ref<!vmla.buffer>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.view(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32, %byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.copy(%src : !vm.ref<!vmla.buffer>, %src_byte_offset : i32, %dst : !vm.ref<!vmla.buffer>, %dst_byte_offset : i32, %byte_length : i32) attributes {sym_visibility = "private"} | |
| vm.import @vmla.buffer.fill(%value : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.buffer.load.i32(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.cmp.i8(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cmp.i16(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cmp.i32(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cmp.f32(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.select.x8(%cond : !vm.ref<!vmla.buffer>, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.select.x16(%cond : !vm.ref<!vmla.buffer>, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.select.x32(%cond : !vm.ref<!vmla.buffer>, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.copy.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %src_indices : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %dst_indices : i32..., %lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.copy.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %src_indices : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %dst_indices : i32..., %lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.copy.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %src_indices : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %dst_indices : i32..., %lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.transpose.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %permutation : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.transpose.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %permutation : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.transpose.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %permutation : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reverse.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dimensions : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reverse.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dimensions : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reverse.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dimensions : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pad.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %value : !vm.ref<!vmla.buffer>, %value_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %edge_padding_low : i32..., %edge_padding_high : i32..., %interior_padding : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pad.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %value : !vm.ref<!vmla.buffer>, %value_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %edge_padding_low : i32..., %edge_padding_high : i32..., %interior_padding : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pad.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %value : !vm.ref<!vmla.buffer>, %value_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %edge_padding_low : i32..., %edge_padding_high : i32..., %interior_padding : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.broadcast.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.broadcast.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.broadcast.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tile.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tile.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tile.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.not.x8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.not.x16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.not.x32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.and.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.and.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.and.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.or.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.or.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.or.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.xor.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.xor.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.xor.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shl.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shl.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shl.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.u8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.u16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.u32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.u8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.u16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.u32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.u8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.u16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.u32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pow.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.exp.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.log.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rsqrt.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sqrt.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cos.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sin.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tanh.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.atan2.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.i8(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.i16(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.i32(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.f32(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.floor.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.ceil.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i8.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i8.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i8.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i16.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i16.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i16.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i32.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i32.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i32.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.f32.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.f32.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.f32.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.conv.f32f32.f32(%input : !vm.ref<!vmla.buffer>, %input_shape : i32..., %filter : !vm.ref<!vmla.buffer>, %filter_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %window_strides : i32..., %padding : i32..., %lhs_dilation : i32..., %rhs_dilation : i32..., %feature_group_count : i32, %batch_group_count : i32) attributes {sym_visibility = "private"} | |
| vm.import @vmla.batch.matmul.f32f32.f32(%lhs : !vm.ref<!vmla.buffer>, %lhs_shape : i32..., %rhs : !vm.ref<!vmla.buffer>, %rhs_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.i8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.i16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.i32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.f32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.i8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.i16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.i32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.f32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.i8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.i16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.i32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.f32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| } | |
| } | |
| *** IR Dump Before SymbolDCE *** | |
| module { | |
| vm.module @module { | |
| vm.func @simple_mul_ex_dispatch_0(%arg0: !vm.ref<!vmla.interface>) { | |
| %c4 = vm.const.i32 4 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %c1 = vm.const.i32 1 : i32 | |
| %0 = vm.call @vmla.interface.const(%arg0, %c1) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %1 = vm.call @vmla.interface.const(%arg0, %zero) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %ref = vm.call @vmla.interface.binding(%arg0, %zero, %zero) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| %2 = vm.mul.i32 %1, %c4 : i32 | |
| %3 = vm.mul.i32 %2, %0 : i32 | |
| %ref_0 = vm.call @vmla.buffer.alloc(%3) : (i32) -> !vm.ref<!vmla.buffer> | |
| %ref_1 = vm.call @vmla.buffer.view(%ref, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.abs.f32(%ref_1, %ref_0) : (!vm.ref<!vmla.buffer>, !vm.ref<!vmla.buffer>) -> () | |
| %ref_2 = vm.call @vmla.interface.binding(%arg0, %zero, %c1) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.buffer.copy(%ref_0, %zero, %ref_2, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, !vm.ref<!vmla.buffer>, i32, i32) -> () | |
| vm.return | |
| } | |
| vm.export @simple_mul_ex_dispatch_0 | |
| vm.import @vmla.interface.current() -> !vm.ref<!vmla.interface> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.interface.const(%interface : !vm.ref<!vmla.interface>, %offset : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.interface.binding(%interface : !vm.ref<!vmla.interface>, %set : i32, %binding : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.const(%value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.alloc(%byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.clone(%src : !vm.ref<!vmla.buffer>) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.byte_length(%value : !vm.ref<!vmla.buffer>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.view(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32, %byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.buffer.copy(%src : !vm.ref<!vmla.buffer>, %src_byte_offset : i32, %dst : !vm.ref<!vmla.buffer>, %dst_byte_offset : i32, %byte_length : i32) attributes {sym_visibility = "private"} | |
| vm.import @vmla.buffer.fill(%value : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.buffer.load.i32(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @vmla.cmp.i8(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cmp.i16(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cmp.i32(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cmp.f32(%predicate : i32, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.select.x8(%cond : !vm.ref<!vmla.buffer>, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.select.x16(%cond : !vm.ref<!vmla.buffer>, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.select.x32(%cond : !vm.ref<!vmla.buffer>, %lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.copy.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %src_indices : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %dst_indices : i32..., %lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.copy.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %src_indices : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %dst_indices : i32..., %lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.copy.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %src_indices : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %dst_indices : i32..., %lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.transpose.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %permutation : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.transpose.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %permutation : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.transpose.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %permutation : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reverse.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dimensions : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reverse.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dimensions : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reverse.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dimensions : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pad.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %value : !vm.ref<!vmla.buffer>, %value_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %edge_padding_low : i32..., %edge_padding_high : i32..., %interior_padding : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pad.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %value : !vm.ref<!vmla.buffer>, %value_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %edge_padding_low : i32..., %edge_padding_high : i32..., %interior_padding : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pad.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %value : !vm.ref<!vmla.buffer>, %value_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %edge_padding_low : i32..., %edge_padding_high : i32..., %interior_padding : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.broadcast.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.broadcast.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.broadcast.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tile.x8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tile.x16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tile.x32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.not.x8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.not.x16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.not.x32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.and.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.and.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.and.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.or.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.or.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.or.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.xor.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.xor.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.xor.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shl.x8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shl.x16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shl.x32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.u8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.u16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.u32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.shr.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.add.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sub.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.abs.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.neg.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.mul.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.u8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.u16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.u32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.div.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.u8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.u16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.u32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rem.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.pow.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.exp.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.log.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.rsqrt.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sqrt.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.cos.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.sin.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.tanh.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.atan2.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.min.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.i8(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.i16(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.i32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.max.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.i8(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.i16(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.i32(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.clamp.f32(%min : !vm.ref<!vmla.buffer>, %value : !vm.ref<!vmla.buffer>, %max : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.floor.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.ceil.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i8.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i8.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i8.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i16.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i16.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i16.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i32.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i32.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.i32.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.f32.i8(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.f32.i16(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.convert.f32.i32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"} | |
| vm.import @vmla.conv.f32f32.f32(%input : !vm.ref<!vmla.buffer>, %input_shape : i32..., %filter : !vm.ref<!vmla.buffer>, %filter_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32..., %window_strides : i32..., %padding : i32..., %lhs_dilation : i32..., %rhs_dilation : i32..., %feature_group_count : i32, %batch_group_count : i32) attributes {sym_visibility = "private"} | |
| vm.import @vmla.batch.matmul.f32f32.f32(%lhs : !vm.ref<!vmla.buffer>, %lhs_shape : i32..., %rhs : !vm.ref<!vmla.buffer>, %rhs_shape : i32..., %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.i8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.i16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.i32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.sum.f32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.i8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.i16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.i32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.min.f32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.i8(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.i16(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.i32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| vm.import @vmla.reduce.max.f32(%src : !vm.ref<!vmla.buffer>, %src_shape : i32..., %init : !vm.ref<!vmla.buffer>, %init_shape : i32..., %dimension : i32, %dst : !vm.ref<!vmla.buffer>, %dst_shape : i32...) attributes {sym_visibility = "private"} | |
| } | |
| } | |
| *** IR Dump Before iree_compiler::(anonymous namespace)::ConvertFlowToHALPass *** | |
| module { | |
| hal.executable @simple_mul_ex_dispatch_0 attributes {sym_visibility = "private"} { | |
| hal.interface @legacy_io attributes {push_constants = 2 : i32} { | |
| hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read" | |
| hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" | |
| } | |
| hal.executable.entry_point @simple_mul_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<?x?xf32>, index, index) -> tensor<?x?xf32>, workgroup_size = [1 : index, 1 : index, 1 : index]} | |
| hal.executable.binary attributes {data = dense<"0x10000000564D4C41000006000800040006000000040000009C03000020000000424D4F44180020001C001800140010000C000000000000000800040018000000D4020000600000005400000048000000280000001400000004000000060000006D6F64756C6500000300000098020000740200005402000006000000FC010000AC010000680100001C010000D40000009C000000010000005800000001000000140000000100000000000000A00000000600040030FEFFFF0800000014000000C2FDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F30000000006CFEFFFF0800000014000000FEFDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000A8FEFFFF08000000180000003AFEFFFF040000000200000001000000010000000C000000766D6C612E6162732E66333200000000DCFEFFFF08000000240000006EFEFFFF0400000005000000010000000000000001000000000000000000000010000000766D6C612E6275666665722E636F70790000000020FFFFFF08000000280000002CFFFFFF080000000C00000001000000010000000300000001000000000000000000000010000000766D6C612E6275666665722E766965770000000068FFFFFF080000002000000074FFFFFF080000000C0000000100000001000000010000000000000011000000766D6C612E6275666665722E616C6C6F63000000A8FFFFFF0800000028000000B4FFFFFF080000000C00000001000000010000000300000002000000000000000000000016000000766D6C612E696E746572666163652E62696E64696E670000F4FFFFFF100000002C00000008000C000800040008000000080000000C000000010000000000000002000000020000000000000014000000766D6C612E696E746572666163652E636F6E737400000000CAFFFFFF040000000F00000021766D6C612E696E7465726661636500E6FFFFFF040000000C00000021766D6C612E627566666572000006000800040006000000040000000300000069333200A0000000090400000000000801000901000000020052000000800200008002000100030052000000800200008001000100040052010000800300008001000100010001802404000000000024000003000000520200008001000000010002805203000080030001C001000000010001805205000080020001C0028000005201000080030000C001000200010000805204000080050002C0010000C0010000000000540000"> : vector<952xi8>, format = 1447906369 : i32} { | |
| vm.module @module { | |
| vm.func @simple_mul_ex_dispatch_0(%arg0: !vm.ref<!vmla.interface>) attributes {ordinal = 0 : i32} { | |
| %c4 = vm.const.i32 4 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %c1 = vm.const.i32 1 : i32 | |
| %0 = vm.call @vmla.interface.const(%arg0, %c1) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %1 = vm.call @vmla.interface.const(%arg0, %zero) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %ref = vm.call @vmla.interface.binding(%arg0, %zero, %zero) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| %2 = vm.mul.i32 %1, %c4 : i32 | |
| %3 = vm.mul.i32 %2, %0 : i32 | |
| %ref_0 = vm.call @vmla.buffer.alloc(%3) : (i32) -> !vm.ref<!vmla.buffer> | |
| %ref_1 = vm.call @vmla.buffer.view(%ref, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.abs.f32(%ref_1, %ref_0) : (!vm.ref<!vmla.buffer>, !vm.ref<!vmla.buffer>) -> () | |
| %ref_2 = vm.call @vmla.interface.binding(%arg0, %zero, %c1) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.buffer.copy(%ref_0, %zero, %ref_2, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, !vm.ref<!vmla.buffer>, i32, i32) -> () | |
| vm.return | |
| } | |
| vm.export @simple_mul_ex_dispatch_0 attributes {ordinal = 0 : i32} | |
| vm.import @vmla.interface.const(%interface : !vm.ref<!vmla.interface>, %offset : i32) -> i32 attributes {nosideeffects, ordinal = 0 : i32, sym_visibility = "private"} | |
| vm.import @vmla.interface.binding(%interface : !vm.ref<!vmla.interface>, %set : i32, %binding : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 1 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.alloc(%byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 2 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.view(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32, %byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 3 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.copy(%src : !vm.ref<!vmla.buffer>, %src_byte_offset : i32, %dst : !vm.ref<!vmla.buffer>, %dst_byte_offset : i32, %byte_length : i32) attributes {ordinal = 4 : i32, sym_visibility = "private"} | |
| vm.import @vmla.abs.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {ordinal = 5 : i32, sym_visibility = "private"} | |
| } | |
| } | |
| } | |
| func @simple_mul(%arg0: tensor<?x?xf32>, %arg1: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %3 = flow.ex.stream.fragment(%arg2 = %arg0 : tensor<?x?xf32>, %arg3 = %0 : index, %arg4 = %1 : index, %arg5 = %2 : index) -> tensor<?x?xf32> { | |
| %4 = shapex.make_ranked_shape %arg3, %arg4 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %5 = shapex.tie_shape %arg2, %4 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| %6 = flow.dispatch @simple_mul_ex_dispatch_0::@simple_mul_ex_dispatch_0[%arg5 : index](%5, %arg3, %arg4) : (tensor<?x?xf32>, index, index) -> tensor<?x?xf32> | |
| %7 = shapex.tie_shape %6, %4 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| flow.return %7 : tensor<?x?xf32> | |
| } | |
| return %3, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?,?]> | |
| } | |
| } | |
| @RESOLVE DIM: %0 = shapex.ranked_dim <<UNKNOWN SSA VALUE>>[0] : !shapex.ranked_shape<[?,?]> -> index | |
| @RESOLVE DIM: %1 = shapex.ranked_dim <<UNKNOWN SSA VALUE>>[1] : !shapex.ranked_shape<[?,?]> -> index | |
| @RESOLVE DIM: %0 = shapex.ranked_dim <<UNKNOWN SSA VALUE>>[0] : !shapex.ranked_shape<[?,?]> -> index | |
| @RESOLVE DIM: %1 = shapex.ranked_dim <<UNKNOWN SSA VALUE>>[1] : !shapex.ranked_shape<[?,?]> -> index | |
| @RESOLVE DIM: %0 = shapex.ranked_dim <<UNKNOWN SSA VALUE>>[0] : !shapex.ranked_shape<[?,?]> -> index | |
| @RESOLVE DIM: %1 = shapex.ranked_dim <<UNKNOWN SSA VALUE>>[1] : !shapex.ranked_shape<[?,?]> -> index | |
| @RESOLVE DIM: %0 = shapex.ranked_dim <<UNKNOWN SSA VALUE>>[0] : !shapex.ranked_shape<[?,?]> -> index | |
| @RESOLVE DIM: %1 = shapex.ranked_dim <<UNKNOWN SSA VALUE>>[1] : !shapex.ranked_shape<[?,?]> -> index | |
| *** IR Dump Before iree_compiler::Shape::(anonymous namespace)::ExpandFunctionRankedShapeDimsPass *** | |
| func @simple_mul(%arg0: !hal.buffer, %arg1: !shapex.ranked_shape<[?,?]>) -> (!hal.buffer, !shapex.ranked_shape<[?,?]>) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = muli %0, %1 : index | |
| %dev = hal.ex.shared_device : !hal.device | |
| %allocator = hal.device.allocator %dev : !hal.allocator | |
| %sz = hal.allocator.compute_size %allocator, shape = [%0, %1], element_type = 50331680 | |
| %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer | |
| hal.ex.defer_release %buffer : !hal.buffer | |
| %sz_0 = hal.allocator.compute_size %allocator, shape = [%0, %1], element_type = 50331680 | |
| %buffer_1 = hal.allocator.allocate %allocator, "DeviceVisible|DeviceLocal", "Transfer|Dispatch", %sz_0 : !hal.buffer | |
| hal.ex.defer_release %buffer_1 : !hal.buffer | |
| %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer | |
| hal.command_buffer.begin %cmd | |
| %exe = hal.executable.lookup %dev, @simple_mul_ex_dispatch_0 : !hal.executable | |
| %executable_layout = hal.executable_layout.lookup %dev, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">]], push_constants = 2 attributes {push_constants = 2 : i32} : !hal.executable_layout | |
| %c1 = constant 1 : index | |
| %c1_2 = constant 1 : index | |
| %3 = addi %2, %c1_2 : index | |
| %4 = subi %3, %c1 : index | |
| %5 = divi_unsigned %4, %c1_2 : index | |
| %6 = muli %5, %4 : index | |
| %7 = subi %2, %6 : index | |
| %8 = cmpi "sge", %7, %c1 : index | |
| %9 = select %8, %7, %c1 : index | |
| %c1_3 = constant 1 : index | |
| %10 = addi %9, %c1_3 : index | |
| %11 = subi %10, %c1 : index | |
| %12 = divi_unsigned %11, %c1_3 : index | |
| %13 = muli %12, %11 : index | |
| %14 = subi %9, %13 : index | |
| %15 = cmpi "sge", %14, %c1 : index | |
| %16 = select %15, %14, %c1 : index | |
| %c1_4 = constant 1 : index | |
| %17 = addi %16, %c1_4 : index | |
| %18 = subi %17, %c1 : index | |
| %19 = divi_unsigned %18, %c1_4 : index | |
| %20 = muli %19, %18 : index | |
| %21 = subi %16, %20 : index | |
| %22 = cmpi "sge", %21, %c1 : index | |
| %23 = select %22, %21, %c1 : index | |
| %24 = index_cast %0 : index to i32 | |
| %25 = index_cast %1 : index to i32 | |
| hal.command_buffer.push_constants %cmd, %executable_layout, offset = 0, values = [%24, %25] : i32 | |
| %c0 = constant 0 : index | |
| %allocator_5 = hal.buffer.allocator %buffer_1 : !hal.allocator | |
| %sz_6 = hal.allocator.compute_size %allocator_5, shape = [%0, %1], element_type = 50331680 | |
| %allocator_7 = hal.buffer.allocator %buffer : !hal.allocator | |
| %sz_8 = hal.allocator.compute_size %allocator_7, shape = [%0, %1], element_type = 50331680 | |
| hal.command_buffer.push_descriptor_set %cmd, %executable_layout, set=0, bindings=[0 = (%buffer_1, %c0, %sz_6), 1 = (%buffer, %c0, %sz_8)] | |
| hal.command_buffer.dispatch %cmd, %exe, entry_point = 0, workgroup_xyz = [%5, %12, %19] | |
| %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32> | |
| hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier] | |
| hal.command_buffer.end %cmd | |
| hal.ex.submit_and_wait %dev, %cmd | |
| return %buffer, %arg1 : !hal.buffer, !shapex.ranked_shape<[?,?]> | |
| } | |
| *** IR Dump Before iree_compiler::IREE::HAL::(anonymous namespace)::PublicABIGenerationPass *** | |
| module { | |
| hal.executable @simple_mul_ex_dispatch_0 attributes {sym_visibility = "private"} { | |
| hal.interface @legacy_io attributes {push_constants = 2 : i32} { | |
| hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read" | |
| hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" | |
| } | |
| hal.executable.entry_point @simple_mul_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<?x?xf32>, index, index) -> tensor<?x?xf32>, workgroup_size = [1 : index, 1 : index, 1 : index]} | |
| hal.executable.binary attributes {data = dense<"0x10000000564D4C41000006000800040006000000040000009C03000020000000424D4F44180020001C001800140010000C000000000000000800040018000000D4020000600000005400000048000000280000001400000004000000060000006D6F64756C6500000300000098020000740200005402000006000000FC010000AC010000680100001C010000D40000009C000000010000005800000001000000140000000100000000000000A00000000600040030FEFFFF0800000014000000C2FDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F30000000006CFEFFFF0800000014000000FEFDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000A8FEFFFF08000000180000003AFEFFFF040000000200000001000000010000000C000000766D6C612E6162732E66333200000000DCFEFFFF08000000240000006EFEFFFF0400000005000000010000000000000001000000000000000000000010000000766D6C612E6275666665722E636F70790000000020FFFFFF08000000280000002CFFFFFF080000000C00000001000000010000000300000001000000000000000000000010000000766D6C612E6275666665722E766965770000000068FFFFFF080000002000000074FFFFFF080000000C0000000100000001000000010000000000000011000000766D6C612E6275666665722E616C6C6F63000000A8FFFFFF0800000028000000B4FFFFFF080000000C00000001000000010000000300000002000000000000000000000016000000766D6C612E696E746572666163652E62696E64696E670000F4FFFFFF100000002C00000008000C000800040008000000080000000C000000010000000000000002000000020000000000000014000000766D6C612E696E746572666163652E636F6E737400000000CAFFFFFF040000000F00000021766D6C612E696E7465726661636500E6FFFFFF040000000C00000021766D6C612E627566666572000006000800040006000000040000000300000069333200A0000000090400000000000801000901000000020052000000800200008002000100030052000000800200008001000100040052010000800300008001000100010001802404000000000024000003000000520200008001000000010002805203000080030001C001000000010001805205000080020001C0028000005201000080030000C001000200010000805204000080050002C0010000C0010000000000540000"> : vector<952xi8>, format = 1447906369 : i32} { | |
| vm.module @module { | |
| vm.func @simple_mul_ex_dispatch_0(%arg0: !vm.ref<!vmla.interface>) attributes {ordinal = 0 : i32} { | |
| %c4 = vm.const.i32 4 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %c1 = vm.const.i32 1 : i32 | |
| %0 = vm.call @vmla.interface.const(%arg0, %c1) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %1 = vm.call @vmla.interface.const(%arg0, %zero) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %ref = vm.call @vmla.interface.binding(%arg0, %zero, %zero) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| %2 = vm.mul.i32 %1, %c4 : i32 | |
| %3 = vm.mul.i32 %2, %0 : i32 | |
| %ref_0 = vm.call @vmla.buffer.alloc(%3) : (i32) -> !vm.ref<!vmla.buffer> | |
| %ref_1 = vm.call @vmla.buffer.view(%ref, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.abs.f32(%ref_1, %ref_0) : (!vm.ref<!vmla.buffer>, !vm.ref<!vmla.buffer>) -> () | |
| %ref_2 = vm.call @vmla.interface.binding(%arg0, %zero, %c1) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.buffer.copy(%ref_0, %zero, %ref_2, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, !vm.ref<!vmla.buffer>, i32, i32) -> () | |
| vm.return | |
| } | |
| vm.export @simple_mul_ex_dispatch_0 attributes {ordinal = 0 : i32} | |
| vm.import @vmla.interface.const(%interface : !vm.ref<!vmla.interface>, %offset : i32) -> i32 attributes {nosideeffects, ordinal = 0 : i32, sym_visibility = "private"} | |
| vm.import @vmla.interface.binding(%interface : !vm.ref<!vmla.interface>, %set : i32, %binding : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 1 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.alloc(%byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 2 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.view(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32, %byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 3 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.copy(%src : !vm.ref<!vmla.buffer>, %src_byte_offset : i32, %dst : !vm.ref<!vmla.buffer>, %dst_byte_offset : i32, %byte_length : i32) attributes {ordinal = 4 : i32, sym_visibility = "private"} | |
| vm.import @vmla.abs.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {ordinal = 5 : i32, sym_visibility = "private"} | |
| } | |
| } | |
| } | |
| func @simple_mul(%arg0: !hal.buffer, %arg1: index, %arg2: index) -> (!hal.buffer, index, index) attributes {iree.module.export, iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %0 = shapex.make_ranked_shape %arg1, %arg2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %1 = shapex.ranked_dim %0[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = shapex.ranked_dim %0[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %3 = muli %1, %2 : index | |
| %dev = hal.ex.shared_device : !hal.device | |
| %allocator = hal.device.allocator %dev : !hal.allocator | |
| %sz = hal.allocator.compute_size %allocator, shape = [%1, %2], element_type = 50331680 | |
| %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer | |
| hal.ex.defer_release %buffer : !hal.buffer | |
| %sz_0 = hal.allocator.compute_size %allocator, shape = [%1, %2], element_type = 50331680 | |
| %buffer_1 = hal.allocator.allocate %allocator, "DeviceVisible|DeviceLocal", "Transfer|Dispatch", %sz_0 : !hal.buffer | |
| hal.ex.defer_release %buffer_1 : !hal.buffer | |
| %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer | |
| hal.command_buffer.begin %cmd | |
| %exe = hal.executable.lookup %dev, @simple_mul_ex_dispatch_0 : !hal.executable | |
| %executable_layout = hal.executable_layout.lookup %dev, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">]], push_constants = 2 attributes {push_constants = 2 : i32} : !hal.executable_layout | |
| %c1 = constant 1 : index | |
| %c1_2 = constant 1 : index | |
| %4 = addi %3, %c1_2 : index | |
| %5 = subi %4, %c1 : index | |
| %6 = divi_unsigned %5, %c1_2 : index | |
| %7 = muli %6, %5 : index | |
| %8 = subi %3, %7 : index | |
| %9 = cmpi "sge", %8, %c1 : index | |
| %10 = select %9, %8, %c1 : index | |
| %c1_3 = constant 1 : index | |
| %11 = addi %10, %c1_3 : index | |
| %12 = subi %11, %c1 : index | |
| %13 = divi_unsigned %12, %c1_3 : index | |
| %14 = muli %13, %12 : index | |
| %15 = subi %10, %14 : index | |
| %16 = cmpi "sge", %15, %c1 : index | |
| %17 = select %16, %15, %c1 : index | |
| %c1_4 = constant 1 : index | |
| %18 = addi %17, %c1_4 : index | |
| %19 = subi %18, %c1 : index | |
| %20 = divi_unsigned %19, %c1_4 : index | |
| %21 = muli %20, %19 : index | |
| %22 = subi %17, %21 : index | |
| %23 = cmpi "sge", %22, %c1 : index | |
| %24 = select %23, %22, %c1 : index | |
| %25 = index_cast %1 : index to i32 | |
| %26 = index_cast %2 : index to i32 | |
| hal.command_buffer.push_constants %cmd, %executable_layout, offset = 0, values = [%25, %26] : i32 | |
| %c0 = constant 0 : index | |
| %allocator_5 = hal.buffer.allocator %buffer_1 : !hal.allocator | |
| %sz_6 = hal.allocator.compute_size %allocator_5, shape = [%1, %2], element_type = 50331680 | |
| %allocator_7 = hal.buffer.allocator %buffer : !hal.allocator | |
| %sz_8 = hal.allocator.compute_size %allocator_7, shape = [%1, %2], element_type = 50331680 | |
| hal.command_buffer.push_descriptor_set %cmd, %executable_layout, set=0, bindings=[0 = (%buffer_1, %c0, %sz_6), 1 = (%buffer, %c0, %sz_8)] | |
| hal.command_buffer.dispatch %cmd, %exe, entry_point = 0, workgroup_xyz = [%6, %13, %20] | |
| %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32> | |
| hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier] | |
| hal.command_buffer.end %cmd | |
| hal.ex.submit_and_wait %dev, %cmd | |
| %27 = shapex.ranked_dim %0[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %28 = shapex.ranked_dim %0[1] : !shapex.ranked_shape<[?,?]> -> index | |
| return %buffer, %27, %28 : !hal.buffer, index, index | |
| } | |
| } | |
| *** IR Dump Before iree_compiler::IREE::HAL::MaterializeResourceCachesPass *** | |
| module { | |
| hal.executable @simple_mul_ex_dispatch_0 attributes {sym_visibility = "private"} { | |
| hal.interface @legacy_io attributes {push_constants = 2 : i32} { | |
| hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read" | |
| hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" | |
| } | |
| hal.executable.entry_point @simple_mul_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<?x?xf32>, index, index) -> tensor<?x?xf32>, workgroup_size = [1 : index, 1 : index, 1 : index]} | |
| hal.executable.binary attributes {data = dense<"0x10000000564D4C41000006000800040006000000040000009C03000020000000424D4F44180020001C001800140010000C000000000000000800040018000000D4020000600000005400000048000000280000001400000004000000060000006D6F64756C6500000300000098020000740200005402000006000000FC010000AC010000680100001C010000D40000009C000000010000005800000001000000140000000100000000000000A00000000600040030FEFFFF0800000014000000C2FDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F30000000006CFEFFFF0800000014000000FEFDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000A8FEFFFF08000000180000003AFEFFFF040000000200000001000000010000000C000000766D6C612E6162732E66333200000000DCFEFFFF08000000240000006EFEFFFF0400000005000000010000000000000001000000000000000000000010000000766D6C612E6275666665722E636F70790000000020FFFFFF08000000280000002CFFFFFF080000000C00000001000000010000000300000001000000000000000000000010000000766D6C612E6275666665722E766965770000000068FFFFFF080000002000000074FFFFFF080000000C0000000100000001000000010000000000000011000000766D6C612E6275666665722E616C6C6F63000000A8FFFFFF0800000028000000B4FFFFFF080000000C00000001000000010000000300000002000000000000000000000016000000766D6C612E696E746572666163652E62696E64696E670000F4FFFFFF100000002C00000008000C000800040008000000080000000C000000010000000000000002000000020000000000000014000000766D6C612E696E746572666163652E636F6E737400000000CAFFFFFF040000000F00000021766D6C612E696E7465726661636500E6FFFFFF040000000C00000021766D6C612E627566666572000006000800040006000000040000000300000069333200A0000000090400000000000801000901000000020052000000800200008002000100030052000000800200008001000100040052010000800300008001000100010001802404000000000024000003000000520200008001000000010002805203000080030001C001000000010001805205000080020001C0028000005201000080030000C001000200010000805204000080050002C0010000C0010000000000540000"> : vector<952xi8>, format = 1447906369 : i32} { | |
| vm.module @module { | |
| vm.func @simple_mul_ex_dispatch_0(%arg0: !vm.ref<!vmla.interface>) attributes {ordinal = 0 : i32} { | |
| %c4 = vm.const.i32 4 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %c1 = vm.const.i32 1 : i32 | |
| %0 = vm.call @vmla.interface.const(%arg0, %c1) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %1 = vm.call @vmla.interface.const(%arg0, %zero) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %ref = vm.call @vmla.interface.binding(%arg0, %zero, %zero) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| %2 = vm.mul.i32 %1, %c4 : i32 | |
| %3 = vm.mul.i32 %2, %0 : i32 | |
| %ref_0 = vm.call @vmla.buffer.alloc(%3) : (i32) -> !vm.ref<!vmla.buffer> | |
| %ref_1 = vm.call @vmla.buffer.view(%ref, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.abs.f32(%ref_1, %ref_0) : (!vm.ref<!vmla.buffer>, !vm.ref<!vmla.buffer>) -> () | |
| %ref_2 = vm.call @vmla.interface.binding(%arg0, %zero, %c1) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.buffer.copy(%ref_0, %zero, %ref_2, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, !vm.ref<!vmla.buffer>, i32, i32) -> () | |
| vm.return | |
| } | |
| vm.export @simple_mul_ex_dispatch_0 attributes {ordinal = 0 : i32} | |
| vm.import @vmla.interface.const(%interface : !vm.ref<!vmla.interface>, %offset : i32) -> i32 attributes {nosideeffects, ordinal = 0 : i32, sym_visibility = "private"} | |
| vm.import @vmla.interface.binding(%interface : !vm.ref<!vmla.interface>, %set : i32, %binding : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 1 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.alloc(%byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 2 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.view(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32, %byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 3 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.copy(%src : !vm.ref<!vmla.buffer>, %src_byte_offset : i32, %dst : !vm.ref<!vmla.buffer>, %dst_byte_offset : i32, %byte_length : i32) attributes {ordinal = 4 : i32, sym_visibility = "private"} | |
| vm.import @vmla.abs.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {ordinal = 5 : i32, sym_visibility = "private"} | |
| } | |
| } | |
| } | |
| func @simple_mul(%arg0: !hal.buffer, %arg1: index, %arg2: index) -> (!hal.buffer, index, index) attributes {iree.module.export = "simple_mul$raw"} { | |
| %0 = shapex.make_ranked_shape %arg1, %arg2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %1 = shapex.ranked_dim %0[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = shapex.ranked_dim %0[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %3 = muli %1, %2 : index | |
| %dev = hal.ex.shared_device : !hal.device | |
| %allocator = hal.device.allocator %dev : !hal.allocator | |
| %sz = hal.allocator.compute_size %allocator, shape = [%1, %2], element_type = 50331680 | |
| %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer | |
| hal.ex.defer_release %buffer : !hal.buffer | |
| %sz_0 = hal.allocator.compute_size %allocator, shape = [%1, %2], element_type = 50331680 | |
| %buffer_1 = hal.allocator.allocate %allocator, "DeviceVisible|DeviceLocal", "Transfer|Dispatch", %sz_0 : !hal.buffer | |
| hal.ex.defer_release %buffer_1 : !hal.buffer | |
| %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer | |
| hal.command_buffer.begin %cmd | |
| %exe = hal.executable.lookup %dev, @simple_mul_ex_dispatch_0 : !hal.executable | |
| %executable_layout = hal.executable_layout.lookup %dev, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">]], push_constants = 2 attributes {push_constants = 2 : i32} : !hal.executable_layout | |
| %c1 = constant 1 : index | |
| %c1_2 = constant 1 : index | |
| %4 = addi %3, %c1_2 : index | |
| %5 = subi %4, %c1 : index | |
| %6 = divi_unsigned %5, %c1_2 : index | |
| %7 = muli %6, %5 : index | |
| %8 = subi %3, %7 : index | |
| %9 = cmpi "sge", %8, %c1 : index | |
| %10 = select %9, %8, %c1 : index | |
| %c1_3 = constant 1 : index | |
| %11 = addi %10, %c1_3 : index | |
| %12 = subi %11, %c1 : index | |
| %13 = divi_unsigned %12, %c1_3 : index | |
| %14 = muli %13, %12 : index | |
| %15 = subi %10, %14 : index | |
| %16 = cmpi "sge", %15, %c1 : index | |
| %17 = select %16, %15, %c1 : index | |
| %c1_4 = constant 1 : index | |
| %18 = addi %17, %c1_4 : index | |
| %19 = subi %18, %c1 : index | |
| %20 = divi_unsigned %19, %c1_4 : index | |
| %21 = muli %20, %19 : index | |
| %22 = subi %17, %21 : index | |
| %23 = cmpi "sge", %22, %c1 : index | |
| %24 = select %23, %22, %c1 : index | |
| %25 = index_cast %1 : index to i32 | |
| %26 = index_cast %2 : index to i32 | |
| hal.command_buffer.push_constants %cmd, %executable_layout, offset = 0, values = [%25, %26] : i32 | |
| %c0 = constant 0 : index | |
| %allocator_5 = hal.buffer.allocator %buffer_1 : !hal.allocator | |
| %sz_6 = hal.allocator.compute_size %allocator_5, shape = [%1, %2], element_type = 50331680 | |
| %allocator_7 = hal.buffer.allocator %buffer : !hal.allocator | |
| %sz_8 = hal.allocator.compute_size %allocator_7, shape = [%1, %2], element_type = 50331680 | |
| hal.command_buffer.push_descriptor_set %cmd, %executable_layout, set=0, bindings=[0 = (%buffer_1, %c0, %sz_6), 1 = (%buffer, %c0, %sz_8)] | |
| hal.command_buffer.dispatch %cmd, %exe, entry_point = 0, workgroup_xyz = [%6, %13, %20] | |
| %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32> | |
| hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier] | |
| hal.command_buffer.end %cmd | |
| hal.ex.submit_and_wait %dev, %cmd | |
| %27 = shapex.ranked_dim %0[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %28 = shapex.ranked_dim %0[1] : !shapex.ranked_shape<[?,?]> -> index | |
| return %buffer, %27, %28 : !hal.buffer, index, index | |
| } | |
| func @simple_mul$sync(%arg0: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.module.export = "simple_mul", iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %buffer = hal.buffer_view.buffer %arg0 : !hal.buffer | |
| %0 = hal.buffer_view.dim %arg0, 0 : index | |
| %1 = hal.buffer_view.dim %arg0, 1 : index | |
| %2:3 = call @simple_mul(%buffer, %0, %1) : (!hal.buffer, index, index) -> (!hal.buffer, index, index) | |
| %view = hal.buffer_view.create %2#0, shape = [%2#1, %2#2], element_type = 50331680 : !hal.buffer_view | |
| return %view : !hal.buffer_view | |
| } | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} { | |
| %dev = hal.ex.shared_device : !hal.device | |
| %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout | |
| return %descriptor_set_layout : !hal.descriptor_set_layout | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} { | |
| %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout | |
| %dev = hal.ex.shared_device : !hal.device | |
| %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 2 : !hal.executable_layout | |
| return %executable_layout : !hal.executable_layout | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| func @simple_mul(%arg0: !hal.buffer, %arg1: index, %arg2: index) -> (!hal.buffer, index, index) attributes {iree.module.export = "simple_mul$raw"} { | |
| %0 = shapex.make_ranked_shape %arg1, %arg2 : (index, index) -> !shapex.ranked_shape<[?,?]> | |
| %1 = shapex.ranked_dim %0[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %2 = shapex.ranked_dim %0[1] : !shapex.ranked_shape<[?,?]> -> index | |
| %3 = muli %1, %2 : index | |
| %dev = hal.ex.shared_device : !hal.device | |
| %allocator = hal.device.allocator %dev : !hal.allocator | |
| %sz = hal.allocator.compute_size %allocator, shape = [%1, %2], element_type = 50331680 | |
| %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer | |
| hal.ex.defer_release %buffer : !hal.buffer | |
| %sz_0 = hal.allocator.compute_size %allocator, shape = [%1, %2], element_type = 50331680 | |
| %buffer_1 = hal.allocator.allocate %allocator, "DeviceVisible|DeviceLocal", "Transfer|Dispatch", %sz_0 : !hal.buffer | |
| hal.ex.defer_release %buffer_1 : !hal.buffer | |
| %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer | |
| hal.command_buffer.begin %cmd | |
| %4 = hal.variable.load @_executable_simple_mul_ex_dispatch_0 : !hal.executable | |
| %5 = hal.variable.load @_executable_layout_0 : !hal.executable_layout | |
| %c1 = constant 1 : index | |
| %c1_2 = constant 1 : index | |
| %6 = addi %3, %c1_2 : index | |
| %7 = subi %6, %c1 : index | |
| %8 = divi_unsigned %7, %c1_2 : index | |
| %9 = muli %8, %7 : index | |
| %10 = subi %3, %9 : index | |
| %11 = cmpi "sge", %10, %c1 : index | |
| %12 = select %11, %10, %c1 : index | |
| %c1_3 = constant 1 : index | |
| %13 = addi %12, %c1_3 : index | |
| %14 = subi %13, %c1 : index | |
| %15 = divi_unsigned %14, %c1_3 : index | |
| %16 = muli %15, %14 : index | |
| %17 = subi %12, %16 : index | |
| %18 = cmpi "sge", %17, %c1 : index | |
| %19 = select %18, %17, %c1 : index | |
| %c1_4 = constant 1 : index | |
| %20 = addi %19, %c1_4 : index | |
| %21 = subi %20, %c1 : index | |
| %22 = divi_unsigned %21, %c1_4 : index | |
| %23 = muli %22, %21 : index | |
| %24 = subi %19, %23 : index | |
| %25 = cmpi "sge", %24, %c1 : index | |
| %26 = select %25, %24, %c1 : index | |
| %27 = index_cast %1 : index to i32 | |
| %28 = index_cast %2 : index to i32 | |
| hal.command_buffer.push_constants %cmd, %5, offset = 0, values = [%27, %28] : i32 | |
| %c0 = constant 0 : index | |
| %allocator_5 = hal.buffer.allocator %buffer_1 : !hal.allocator | |
| %sz_6 = hal.allocator.compute_size %allocator_5, shape = [%1, %2], element_type = 50331680 | |
| %allocator_7 = hal.buffer.allocator %buffer : !hal.allocator | |
| %sz_8 = hal.allocator.compute_size %allocator_7, shape = [%1, %2], element_type = 50331680 | |
| hal.command_buffer.push_descriptor_set %cmd, %5, set=0, bindings=[0 = (%buffer_1, %c0, %sz_6), 1 = (%buffer, %c0, %sz_8)] | |
| hal.command_buffer.dispatch %cmd, %4, entry_point = 0, workgroup_xyz = [%8, %15, %22] | |
| %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32> | |
| hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier] | |
| hal.command_buffer.end %cmd | |
| hal.ex.submit_and_wait %dev, %cmd | |
| %29 = shapex.ranked_dim %0[0] : !shapex.ranked_shape<[?,?]> -> index | |
| %30 = shapex.ranked_dim %0[1] : !shapex.ranked_shape<[?,?]> -> index | |
| return %buffer, %29, %30 : !hal.buffer, index, index | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} { | |
| %dev = hal.ex.shared_device : !hal.device | |
| %executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache | |
| %0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout | |
| %executable_simple_mul_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @simple_mul_ex_dispatch_0 : !hal.executable | |
| hal.variable.store %executable_simple_mul_ex_dispatch_0, @_executable_simple_mul_ex_dispatch_0 : !hal.executable | |
| return %executable_cache_default : !hal.executable_cache | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| func @simple_mul$sync(%arg0: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.module.export = "simple_mul", iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %buffer = hal.buffer_view.buffer %arg0 : !hal.buffer | |
| %0 = hal.buffer_view.dim %arg0, 0 : index | |
| %1 = hal.buffer_view.dim %arg0, 1 : index | |
| %2:3 = call @simple_mul(%buffer, %0, %1) : (!hal.buffer, index, index) -> (!hal.buffer, index, index) | |
| %view = hal.buffer_view.create %2#0, shape = [%2#1, %2#2], element_type = 50331680 : !hal.buffer_view | |
| return %view : !hal.buffer_view | |
| } | |
| *** IR Dump Before CSE *** | |
| func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} { | |
| %dev = hal.ex.shared_device : !hal.device | |
| %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout | |
| return %descriptor_set_layout : !hal.descriptor_set_layout | |
| } | |
| *** IR Dump Before CSE *** | |
| func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} { | |
| %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout | |
| %dev = hal.ex.shared_device : !hal.device | |
| %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 2 : !hal.executable_layout | |
| return %executable_layout : !hal.executable_layout | |
| } | |
| *** IR Dump Before CSE *** | |
| func @simple_mul(%arg0: !hal.buffer, %arg1: index, %arg2: index) -> (!hal.buffer, index, index) attributes {iree.module.export = "simple_mul$raw"} { | |
| %c1 = constant 1 : index | |
| %c0 = constant 0 : index | |
| %0 = muli %arg1, %arg2 : index | |
| %dev = hal.ex.shared_device : !hal.device | |
| %allocator = hal.device.allocator %dev : !hal.allocator | |
| %sz = hal.allocator.compute_size %allocator, shape = [%arg1, %arg2], element_type = 50331680 | |
| %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer | |
| hal.ex.defer_release %buffer : !hal.buffer | |
| %sz_0 = hal.allocator.compute_size %allocator, shape = [%arg1, %arg2], element_type = 50331680 | |
| %buffer_1 = hal.allocator.allocate %allocator, "DeviceVisible|DeviceLocal", "Transfer|Dispatch", %sz_0 : !hal.buffer | |
| hal.ex.defer_release %buffer_1 : !hal.buffer | |
| %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer | |
| hal.command_buffer.begin %cmd | |
| %1 = hal.variable.load @_executable_simple_mul_ex_dispatch_0 : !hal.executable | |
| %2 = hal.variable.load @_executable_layout_0 : !hal.executable_layout | |
| %3 = addi %0, %c1 : index | |
| %4 = subi %3, %c1 : index | |
| %5 = divi_unsigned %4, %c1 : index | |
| %6 = muli %5, %4 : index | |
| %7 = subi %0, %6 : index | |
| %8 = cmpi "sge", %7, %c1 : index | |
| %9 = select %8, %7, %c1 : index | |
| %10 = addi %9, %c1 : index | |
| %11 = subi %10, %c1 : index | |
| %12 = divi_unsigned %11, %c1 : index | |
| %13 = muli %12, %11 : index | |
| %14 = subi %9, %13 : index | |
| %15 = cmpi "sge", %14, %c1 : index | |
| %16 = select %15, %14, %c1 : index | |
| %17 = addi %16, %c1 : index | |
| %18 = subi %17, %c1 : index | |
| %19 = divi_unsigned %18, %c1 : index | |
| %20 = index_cast %arg1 : index to i32 | |
| %21 = index_cast %arg2 : index to i32 | |
| hal.command_buffer.push_constants %cmd, %2, offset = 0, values = [%20, %21] : i32 | |
| %sz_2 = hal.allocator.compute_size %allocator, shape = [%arg1, %arg2], element_type = 50331680 | |
| %sz_3 = hal.allocator.compute_size %allocator, shape = [%arg1, %arg2], element_type = 50331680 | |
| hal.command_buffer.push_descriptor_set %cmd, %2, set=0, bindings=[0 = (%buffer_1, %c0, %sz_2), 1 = (%buffer, %c0, %sz_3)] | |
| hal.command_buffer.dispatch %cmd, %1, entry_point = 0, workgroup_xyz = [%5, %12, %19] | |
| %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32> | |
| hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier] | |
| hal.command_buffer.end %cmd | |
| hal.ex.submit_and_wait %dev, %cmd | |
| return %buffer, %arg1, %arg2 : !hal.buffer, index, index | |
| } | |
| *** IR Dump Before CSE *** | |
| func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} { | |
| %dev = hal.ex.shared_device : !hal.device | |
| %executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache | |
| %0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout | |
| %executable_simple_mul_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @simple_mul_ex_dispatch_0 : !hal.executable | |
| hal.variable.store %executable_simple_mul_ex_dispatch_0, @_executable_simple_mul_ex_dispatch_0 : !hal.executable | |
| return %executable_cache_default : !hal.executable_cache | |
| } | |
| *** IR Dump Before CSE *** | |
| func @simple_mul$sync(%arg0: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.module.export = "simple_mul", iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %buffer = hal.buffer_view.buffer %arg0 : !hal.buffer | |
| %0 = hal.buffer_view.dim %arg0, 0 : index | |
| %1 = hal.buffer_view.dim %arg0, 1 : index | |
| %2:3 = call @simple_mul(%buffer, %0, %1) : (!hal.buffer, index, index) -> (!hal.buffer, index, index) | |
| %view = hal.buffer_view.create %2#0, shape = [%2#1, %2#2], element_type = 50331680 : !hal.buffer_view | |
| return %view : !hal.buffer_view | |
| } | |
| *** IR Dump Before Canonicalizer *** | |
| module { | |
| hal.variable @_executable_simple_mul_ex_dispatch_0 mutable : !hal.executable | |
| hal.variable @_descriptor_set_layout_0 init(@_descriptor_set_layout_0_initializer) : !hal.descriptor_set_layout | |
| func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} { | |
| %dev = hal.ex.shared_device : !hal.device | |
| %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout | |
| return %descriptor_set_layout : !hal.descriptor_set_layout | |
| } | |
| hal.variable @_executable_layout_0 init(@_executable_layout_0_initializer) : !hal.executable_layout | |
| func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} { | |
| %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout | |
| %dev = hal.ex.shared_device : !hal.device | |
| %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 2 : !hal.executable_layout | |
| return %executable_layout : !hal.executable_layout | |
| } | |
| hal.variable @_executable_cache init(@_executable_cache_initializer) : !hal.executable_cache | |
| func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} { | |
| %dev = hal.ex.shared_device : !hal.device | |
| %executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache | |
| %0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout | |
| %executable_simple_mul_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @simple_mul_ex_dispatch_0 : !hal.executable | |
| hal.variable.store %executable_simple_mul_ex_dispatch_0, @_executable_simple_mul_ex_dispatch_0 : !hal.executable | |
| return %executable_cache_default : !hal.executable_cache | |
| } | |
| hal.executable @simple_mul_ex_dispatch_0 attributes {sym_visibility = "private"} { | |
| hal.interface @legacy_io attributes {push_constants = 2 : i32} { | |
| hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read" | |
| hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" | |
| } | |
| hal.executable.entry_point @simple_mul_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<?x?xf32>, index, index) -> tensor<?x?xf32>, workgroup_size = [1 : index, 1 : index, 1 : index]} | |
| hal.executable.binary attributes {data = dense<"0x10000000564D4C41000006000800040006000000040000009C03000020000000424D4F44180020001C001800140010000C000000000000000800040018000000D4020000600000005400000048000000280000001400000004000000060000006D6F64756C6500000300000098020000740200005402000006000000FC010000AC010000680100001C010000D40000009C000000010000005800000001000000140000000100000000000000A00000000600040030FEFFFF0800000014000000C2FDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F30000000006CFEFFFF0800000014000000FEFDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000A8FEFFFF08000000180000003AFEFFFF040000000200000001000000010000000C000000766D6C612E6162732E66333200000000DCFEFFFF08000000240000006EFEFFFF0400000005000000010000000000000001000000000000000000000010000000766D6C612E6275666665722E636F70790000000020FFFFFF08000000280000002CFFFFFF080000000C00000001000000010000000300000001000000000000000000000010000000766D6C612E6275666665722E766965770000000068FFFFFF080000002000000074FFFFFF080000000C0000000100000001000000010000000000000011000000766D6C612E6275666665722E616C6C6F63000000A8FFFFFF0800000028000000B4FFFFFF080000000C00000001000000010000000300000002000000000000000000000016000000766D6C612E696E746572666163652E62696E64696E670000F4FFFFFF100000002C00000008000C000800040008000000080000000C000000010000000000000002000000020000000000000014000000766D6C612E696E746572666163652E636F6E737400000000CAFFFFFF040000000F00000021766D6C612E696E7465726661636500E6FFFFFF040000000C00000021766D6C612E627566666572000006000800040006000000040000000300000069333200A0000000090400000000000801000901000000020052000000800200008002000100030052000000800200008001000100040052010000800300008001000100010001802404000000000024000003000000520200008001000000010002805203000080030001C001000000010001805205000080020001C0028000005201000080030000C001000200010000805204000080050002C0010000C0010000000000540000"> : vector<952xi8>, format = 1447906369 : i32} { | |
| vm.module @module { | |
| vm.func @simple_mul_ex_dispatch_0(%arg0: !vm.ref<!vmla.interface>) attributes {ordinal = 0 : i32} { | |
| %c4 = vm.const.i32 4 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %c1 = vm.const.i32 1 : i32 | |
| %0 = vm.call @vmla.interface.const(%arg0, %c1) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %1 = vm.call @vmla.interface.const(%arg0, %zero) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %ref = vm.call @vmla.interface.binding(%arg0, %zero, %zero) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| %2 = vm.mul.i32 %1, %c4 : i32 | |
| %3 = vm.mul.i32 %2, %0 : i32 | |
| %ref_0 = vm.call @vmla.buffer.alloc(%3) : (i32) -> !vm.ref<!vmla.buffer> | |
| %ref_1 = vm.call @vmla.buffer.view(%ref, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.abs.f32(%ref_1, %ref_0) : (!vm.ref<!vmla.buffer>, !vm.ref<!vmla.buffer>) -> () | |
| %ref_2 = vm.call @vmla.interface.binding(%arg0, %zero, %c1) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.buffer.copy(%ref_0, %zero, %ref_2, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, !vm.ref<!vmla.buffer>, i32, i32) -> () | |
| vm.return | |
| } | |
| vm.export @simple_mul_ex_dispatch_0 attributes {ordinal = 0 : i32} | |
| vm.import @vmla.interface.const(%interface : !vm.ref<!vmla.interface>, %offset : i32) -> i32 attributes {nosideeffects, ordinal = 0 : i32, sym_visibility = "private"} | |
| vm.import @vmla.interface.binding(%interface : !vm.ref<!vmla.interface>, %set : i32, %binding : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 1 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.alloc(%byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 2 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.view(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32, %byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 3 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.copy(%src : !vm.ref<!vmla.buffer>, %src_byte_offset : i32, %dst : !vm.ref<!vmla.buffer>, %dst_byte_offset : i32, %byte_length : i32) attributes {ordinal = 4 : i32, sym_visibility = "private"} | |
| vm.import @vmla.abs.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {ordinal = 5 : i32, sym_visibility = "private"} | |
| } | |
| } | |
| } | |
| func @simple_mul(%arg0: !hal.buffer, %arg1: index, %arg2: index) -> (!hal.buffer, index, index) attributes {iree.module.export = "simple_mul$raw"} { | |
| %c1 = constant 1 : index | |
| %c0 = constant 0 : index | |
| %0 = muli %arg1, %arg2 : index | |
| %dev = hal.ex.shared_device : !hal.device | |
| %allocator = hal.device.allocator %dev : !hal.allocator | |
| %sz = hal.allocator.compute_size %allocator, shape = [%arg1, %arg2], element_type = 50331680 | |
| %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer | |
| hal.ex.defer_release %buffer : !hal.buffer | |
| %buffer_0 = hal.allocator.allocate %allocator, "DeviceVisible|DeviceLocal", "Transfer|Dispatch", %sz : !hal.buffer | |
| hal.ex.defer_release %buffer_0 : !hal.buffer | |
| %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer | |
| hal.command_buffer.begin %cmd | |
| %1 = hal.variable.load @_executable_simple_mul_ex_dispatch_0 : !hal.executable | |
| %2 = hal.variable.load @_executable_layout_0 : !hal.executable_layout | |
| %3 = addi %0, %c1 : index | |
| %4 = subi %3, %c1 : index | |
| %5 = divi_unsigned %4, %c1 : index | |
| %6 = muli %5, %4 : index | |
| %7 = subi %0, %6 : index | |
| %8 = cmpi "sge", %7, %c1 : index | |
| %9 = select %8, %7, %c1 : index | |
| %10 = addi %9, %c1 : index | |
| %11 = subi %10, %c1 : index | |
| %12 = divi_unsigned %11, %c1 : index | |
| %13 = muli %12, %11 : index | |
| %14 = subi %9, %13 : index | |
| %15 = cmpi "sge", %14, %c1 : index | |
| %16 = select %15, %14, %c1 : index | |
| %17 = addi %16, %c1 : index | |
| %18 = subi %17, %c1 : index | |
| %19 = divi_unsigned %18, %c1 : index | |
| %20 = index_cast %arg1 : index to i32 | |
| %21 = index_cast %arg2 : index to i32 | |
| hal.command_buffer.push_constants %cmd, %2, offset = 0, values = [%20, %21] : i32 | |
| hal.command_buffer.push_descriptor_set %cmd, %2, set=0, bindings=[0 = (%buffer_0, %c0, %sz), 1 = (%buffer, %c0, %sz)] | |
| hal.command_buffer.dispatch %cmd, %1, entry_point = 0, workgroup_xyz = [%5, %12, %19] | |
| %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32> | |
| hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier] | |
| hal.command_buffer.end %cmd | |
| hal.ex.submit_and_wait %dev, %cmd | |
| return %buffer, %arg1, %arg2 : !hal.buffer, index, index | |
| } | |
| func @simple_mul$sync(%arg0: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.module.export = "simple_mul", iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %buffer = hal.buffer_view.buffer %arg0 : !hal.buffer | |
| %0 = hal.buffer_view.dim %arg0, 0 : index | |
| %1 = hal.buffer_view.dim %arg0, 1 : index | |
| %2:3 = call @simple_mul(%buffer, %0, %1) : (!hal.buffer, index, index) -> (!hal.buffer, index, index) | |
| %view = hal.buffer_view.create %2#0, shape = [%2#1, %2#2], element_type = 50331680 : !hal.buffer_view | |
| return %view : !hal.buffer_view | |
| } | |
| } | |
| *** IR Dump Before iree_compiler::IREE::VM::ConversionPass *** | |
| module { | |
| hal.variable @_executable_simple_mul_ex_dispatch_0 mutable : !hal.executable | |
| hal.variable @_descriptor_set_layout_0 init(@_descriptor_set_layout_0_initializer) : !hal.descriptor_set_layout | |
| func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} { | |
| %dev = hal.ex.shared_device : !hal.device | |
| %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout | |
| return %descriptor_set_layout : !hal.descriptor_set_layout | |
| } | |
| hal.variable @_executable_layout_0 init(@_executable_layout_0_initializer) : !hal.executable_layout | |
| func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} { | |
| %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout | |
| %dev = hal.ex.shared_device : !hal.device | |
| %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 2 : !hal.executable_layout | |
| return %executable_layout : !hal.executable_layout | |
| } | |
| hal.variable @_executable_cache init(@_executable_cache_initializer) : !hal.executable_cache | |
| func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} { | |
| %dev = hal.ex.shared_device : !hal.device | |
| %executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache | |
| %0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout | |
| %executable_simple_mul_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @simple_mul_ex_dispatch_0 : !hal.executable | |
| hal.variable.store %executable_simple_mul_ex_dispatch_0, @_executable_simple_mul_ex_dispatch_0 : !hal.executable | |
| return %executable_cache_default : !hal.executable_cache | |
| } | |
| hal.executable @simple_mul_ex_dispatch_0 attributes {sym_visibility = "private"} { | |
| hal.interface @legacy_io attributes {push_constants = 2 : i32} { | |
| hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read" | |
| hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" | |
| } | |
| hal.executable.entry_point @simple_mul_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<?x?xf32>, index, index) -> tensor<?x?xf32>, workgroup_size = [1 : index, 1 : index, 1 : index]} | |
| hal.executable.binary attributes {data = dense<"0x10000000564D4C41000006000800040006000000040000009C03000020000000424D4F44180020001C001800140010000C000000000000000800040018000000D4020000600000005400000048000000280000001400000004000000060000006D6F64756C6500000300000098020000740200005402000006000000FC010000AC010000680100001C010000D40000009C000000010000005800000001000000140000000100000000000000A00000000600040030FEFFFF0800000014000000C2FDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F30000000006CFEFFFF0800000014000000FEFDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000A8FEFFFF08000000180000003AFEFFFF040000000200000001000000010000000C000000766D6C612E6162732E66333200000000DCFEFFFF08000000240000006EFEFFFF0400000005000000010000000000000001000000000000000000000010000000766D6C612E6275666665722E636F70790000000020FFFFFF08000000280000002CFFFFFF080000000C00000001000000010000000300000001000000000000000000000010000000766D6C612E6275666665722E766965770000000068FFFFFF080000002000000074FFFFFF080000000C0000000100000001000000010000000000000011000000766D6C612E6275666665722E616C6C6F63000000A8FFFFFF0800000028000000B4FFFFFF080000000C00000001000000010000000300000002000000000000000000000016000000766D6C612E696E746572666163652E62696E64696E670000F4FFFFFF100000002C00000008000C000800040008000000080000000C000000010000000000000002000000020000000000000014000000766D6C612E696E746572666163652E636F6E737400000000CAFFFFFF040000000F00000021766D6C612E696E7465726661636500E6FFFFFF040000000C00000021766D6C612E627566666572000006000800040006000000040000000300000069333200A0000000090400000000000801000901000000020052000000800200008002000100030052000000800200008001000100040052010000800300008001000100010001802404000000000024000003000000520200008001000000010002805203000080030001C001000000010001805205000080020001C0028000005201000080030000C001000200010000805204000080050002C0010000C0010000000000540000"> : vector<952xi8>, format = 1447906369 : i32} { | |
| vm.module @module { | |
| vm.func @simple_mul_ex_dispatch_0(%arg0: !vm.ref<!vmla.interface>) attributes {ordinal = 0 : i32} { | |
| %c4 = vm.const.i32 4 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %c1 = vm.const.i32 1 : i32 | |
| %0 = vm.call @vmla.interface.const(%arg0, %c1) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %1 = vm.call @vmla.interface.const(%arg0, %zero) : (!vm.ref<!vmla.interface>, i32) -> i32 | |
| %ref = vm.call @vmla.interface.binding(%arg0, %zero, %zero) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| %2 = vm.mul.i32 %1, %c4 : i32 | |
| %3 = vm.mul.i32 %2, %0 : i32 | |
| %ref_0 = vm.call @vmla.buffer.alloc(%3) : (i32) -> !vm.ref<!vmla.buffer> | |
| %ref_1 = vm.call @vmla.buffer.view(%ref, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.abs.f32(%ref_1, %ref_0) : (!vm.ref<!vmla.buffer>, !vm.ref<!vmla.buffer>) -> () | |
| %ref_2 = vm.call @vmla.interface.binding(%arg0, %zero, %c1) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer> | |
| vm.call @vmla.buffer.copy(%ref_0, %zero, %ref_2, %zero, %3) : (!vm.ref<!vmla.buffer>, i32, !vm.ref<!vmla.buffer>, i32, i32) -> () | |
| vm.return | |
| } | |
| vm.export @simple_mul_ex_dispatch_0 attributes {ordinal = 0 : i32} | |
| vm.import @vmla.interface.const(%interface : !vm.ref<!vmla.interface>, %offset : i32) -> i32 attributes {nosideeffects, ordinal = 0 : i32, sym_visibility = "private"} | |
| vm.import @vmla.interface.binding(%interface : !vm.ref<!vmla.interface>, %set : i32, %binding : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 1 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.alloc(%byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 2 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.view(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32, %byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, ordinal = 3 : i32, sym_visibility = "private"} | |
| vm.import @vmla.buffer.copy(%src : !vm.ref<!vmla.buffer>, %src_byte_offset : i32, %dst : !vm.ref<!vmla.buffer>, %dst_byte_offset : i32, %byte_length : i32) attributes {ordinal = 4 : i32, sym_visibility = "private"} | |
| vm.import @vmla.abs.f32(%src : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {ordinal = 5 : i32, sym_visibility = "private"} | |
| } | |
| } | |
| } | |
| func @simple_mul(%arg0: !hal.buffer, %arg1: index, %arg2: index) -> (!hal.buffer, index, index) attributes {iree.module.export = "simple_mul$raw"} { | |
| %c1 = constant 1 : index | |
| %c0 = constant 0 : index | |
| %0 = muli %arg1, %arg2 : index | |
| %dev = hal.ex.shared_device : !hal.device | |
| %allocator = hal.device.allocator %dev : !hal.allocator | |
| %sz = hal.allocator.compute_size %allocator, shape = [%arg1, %arg2], element_type = 50331680 | |
| %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer | |
| hal.ex.defer_release %buffer : !hal.buffer | |
| %buffer_0 = hal.allocator.allocate %allocator, "DeviceVisible|DeviceLocal", "Transfer|Dispatch", %sz : !hal.buffer | |
| hal.ex.defer_release %buffer_0 : !hal.buffer | |
| %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer | |
| hal.command_buffer.begin %cmd | |
| %1 = hal.variable.load @_executable_simple_mul_ex_dispatch_0 : !hal.executable | |
| %2 = hal.variable.load @_executable_layout_0 : !hal.executable_layout | |
| %3 = addi %0, %c1 : index | |
| %4 = subi %3, %c1 : index | |
| %5 = divi_unsigned %4, %c1 : index | |
| %6 = muli %5, %4 : index | |
| %7 = subi %0, %6 : index | |
| %8 = cmpi "sge", %7, %c1 : index | |
| %9 = select %8, %7, %c1 : index | |
| %10 = addi %9, %c1 : index | |
| %11 = subi %10, %c1 : index | |
| %12 = divi_unsigned %11, %c1 : index | |
| %13 = muli %12, %11 : index | |
| %14 = subi %9, %13 : index | |
| %15 = cmpi "sge", %14, %c1 : index | |
| %16 = select %15, %14, %c1 : index | |
| %17 = addi %16, %c1 : index | |
| %18 = subi %17, %c1 : index | |
| %19 = divi_unsigned %18, %c1 : index | |
| %20 = index_cast %arg1 : index to i32 | |
| %21 = index_cast %arg2 : index to i32 | |
| hal.command_buffer.push_constants %cmd, %2, offset = 0, values = [%20, %21] : i32 | |
| hal.command_buffer.push_descriptor_set %cmd, %2, set=0, bindings=[0 = (%buffer_0, %c0, %sz), 1 = (%buffer, %c0, %sz)] | |
| hal.command_buffer.dispatch %cmd, %1, entry_point = 0, workgroup_xyz = [%5, %12, %19] | |
| %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32> | |
| hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier] | |
| hal.command_buffer.end %cmd | |
| hal.ex.submit_and_wait %dev, %cmd | |
| return %buffer, %arg1, %arg2 : !hal.buffer, index, index | |
| } | |
| func @simple_mul$sync(%arg0: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.module.export = "simple_mul", iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %buffer = hal.buffer_view.buffer %arg0 : !hal.buffer | |
| %0 = hal.buffer_view.dim %arg0, 0 : index | |
| %1 = hal.buffer_view.dim %arg0, 1 : index | |
| %2:3 = call @simple_mul(%buffer, %0, %1) : (!hal.buffer, index, index) -> (!hal.buffer, index, index) | |
| %view = hal.buffer_view.create %2#0, shape = [%2#1, %2#2], element_type = 50331680 : !hal.buffer_view | |
| return %view : !hal.buffer_view | |
| } | |
| } | |
| *** IR Dump Before iree_compiler::IREE::VM::GlobalInitializationPass *** | |
| vm.module @module { | |
| vm.global.ref @_executable_simple_mul_ex_dispatch_0 mutable : !vm.ref<!hal.executable> | |
| vm.global.ref @_descriptor_set_layout_0 init(@_descriptor_set_layout_0_initializer) : !vm.ref<!hal.descriptor_set_layout> attributes {initializer = @_descriptor_set_layout_0_initializer} | |
| vm.func @_descriptor_set_layout_0_initializer() -> !vm.ref<!hal.descriptor_set_layout> attributes {sym_visibility = "private"} { | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %c1 = vm.const.i32 1 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %c7 = vm.const.i32 7 : i32 | |
| %c1_0 = vm.const.i32 1 : i32 | |
| %c1_1 = vm.const.i32 1 : i32 | |
| %c7_2 = vm.const.i32 7 : i32 | |
| %c6 = vm.const.i32 6 : i32 | |
| %ref_3 = vm.call.variadic @hal.descriptor_set_layout.create(%ref, %c1, [%zero, %c7, %c1_0, %c1_1, %c7_2, %c6]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> | |
| vm.return %ref_3 : !vm.ref<!hal.descriptor_set_layout> | |
| } | |
| vm.global.ref @_executable_layout_0 init(@_executable_layout_0_initializer) : !vm.ref<!hal.executable_layout> attributes {initializer = @_executable_layout_0_initializer} | |
| vm.func @_executable_layout_0_initializer() -> !vm.ref<!hal.executable_layout> attributes {sym_visibility = "private"} { | |
| %_descriptor_set_layout_0 = vm.global.load.ref @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout> | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %c2 = vm.const.i32 2 : i32 | |
| %ref_0 = vm.call.variadic @hal.executable_layout.create(%ref, [%_descriptor_set_layout_0], %c2) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout> | |
| vm.return %ref_0 : !vm.ref<!hal.executable_layout> | |
| } | |
| vm.global.ref @_executable_cache init(@_executable_cache_initializer) : !vm.ref<!hal.executable_cache> attributes {initializer = @_executable_cache_initializer} | |
| vm.rodata @_utf8_default_7FD5254DFCA3A5D0 dense<[100, 101, 102, 97, 117, 108, 116]> : vector<7xi8> | |
| vm.rodata @_simple_mul_ex_dispatch_0_binary_vmla dense<"0x10000000564D4C41000006000800040006000000040000009C03000020000000424D4F44180020001C001800140010000C000000000000000800040018000000D4020000600000005400000048000000280000001400000004000000060000006D6F64756C6500000300000098020000740200005402000006000000FC010000AC010000680100001C010000D40000009C000000010000005800000001000000140000000100000000000000A00000000600040030FEFFFF0800000014000000C2FDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F30000000006CFEFFFF0800000014000000FEFDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000A8FEFFFF08000000180000003AFEFFFF040000000200000001000000010000000C000000766D6C612E6162732E66333200000000DCFEFFFF08000000240000006EFEFFFF0400000005000000010000000000000001000000000000000000000010000000766D6C612E6275666665722E636F70790000000020FFFFFF08000000280000002CFFFFFF080000000C00000001000000010000000300000001000000000000000000000010000000766D6C612E6275666665722E766965770000000068FFFFFF080000002000000074FFFFFF080000000C0000000100000001000000010000000000000011000000766D6C612E6275666665722E616C6C6F63000000A8FFFFFF0800000028000000B4FFFFFF080000000C00000001000000010000000300000002000000000000000000000016000000766D6C612E696E746572666163652E62696E64696E670000F4FFFFFF100000002C00000008000C000800040008000000080000000C000000010000000000000002000000020000000000000014000000766D6C612E696E746572666163652E636F6E737400000000CAFFFFFF040000000F00000021766D6C612E696E7465726661636500E6FFFFFF040000000C00000021766D6C612E627566666572000006000800040006000000040000000300000069333200A0000000090400000000000801000901000000020052000000800200008002000100030052000000800200008001000100040052010000800300008001000100010001802404000000000024000003000000520200008001000000010002805203000080030001C001000000010001805205000080020001C0028000005201000080030000C001000200010000805204000080050002C0010000C0010000000000540000"> : vector<952xi8> | |
| vm.func @_executable_cache_initializer() -> !vm.ref<!hal.executable_cache> attributes {sym_visibility = "private"} { | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %_utf8_default_7FD5254DFCA3A5D0 = vm.const.ref.rodata @_utf8_default_7FD5254DFCA3A5D0 : !vm.ref<!iree.byte_buffer> | |
| %ref_0 = vm.call @hal.executable_cache.create(%ref, %_utf8_default_7FD5254DFCA3A5D0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> | |
| %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout> | |
| %c1447906369 = vm.const.i32 1447906369 : i32 | |
| %0 = vm.call.variadic @hal.executable_cache.select_format(%ref_0, [%c1447906369]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32 | |
| %_simple_mul_ex_dispatch_0_binary_vmla = vm.const.ref.rodata @_simple_mul_ex_dispatch_0_binary_vmla : !vm.ref<!iree.byte_buffer> | |
| %null = vm.const.ref.zero : !vm.ref<!iree.byte_buffer> | |
| %ref_1 = vm.switch.ref %0[%_simple_mul_ex_dispatch_0_binary_vmla] else %null : !vm.ref<!iree.byte_buffer> | |
| %c7 = vm.const.i32 7 : i32 | |
| %ref_2 = vm.call @hal.executable_cache.prepare(%ref_0, %_executable_layout_0, %c7, %ref_1) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> | |
| vm.global.store.ref %ref_2, @_executable_simple_mul_ex_dispatch_0 : !vm.ref<!hal.executable> | |
| vm.return %ref_0 : !vm.ref<!hal.executable_cache> | |
| } | |
| vm.func @simple_mul(%arg0: !vm.ref<!hal.buffer>, %arg1: i32, %arg2: i32) -> (!vm.ref<!hal.buffer>, i32, i32) { | |
| %c1 = vm.const.i32 1 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %0 = vm.mul.i32 %arg1, %arg2 : i32 | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %ref_0 = vm.call @hal.device.allocator(%ref) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> | |
| %c50331680 = vm.const.i32 50331680 : i32 | |
| %1 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [%arg1, %arg2], %c50331680) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32 | |
| %c50 = vm.const.i32 50 : i32 | |
| %c15 = vm.const.i32 15 : i32 | |
| %ref_1 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %1) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer> | |
| vm.call @hal.ex.defer_release(%ref_1) : (!vm.ref<!hal.buffer>) -> () | |
| %c48 = vm.const.i32 48 : i32 | |
| %c10 = vm.const.i32 10 : i32 | |
| %ref_2 = vm.call @hal.allocator.allocate(%ref_0, %c48, %c10, %1) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer> | |
| vm.call @hal.ex.defer_release(%ref_2) : (!vm.ref<!hal.buffer>) -> () | |
| %c1_3 = vm.const.i32 1 : i32 | |
| %c3 = vm.const.i32 3 : i32 | |
| %ref_4 = vm.call @hal.command_buffer.create(%ref, %c1_3, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer> | |
| vm.call @hal.command_buffer.begin(%ref_4) : (!vm.ref<!hal.command_buffer>) -> () | |
| %_executable_simple_mul_ex_dispatch_0 = vm.global.load.ref @_executable_simple_mul_ex_dispatch_0 : !vm.ref<!hal.executable> | |
| %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout> | |
| %2 = vm.add.i32 %0, %c1 : i32 | |
| %3 = vm.sub.i32 %2, %c1 : i32 | |
| %4 = vm.div.i32.u %3, %c1 : i32 | |
| %5 = vm.mul.i32 %4, %3 : i32 | |
| %6 = vm.sub.i32 %0, %5 : i32 | |
| %sgte = vm.cmp.gte.i32.s %6, %c1 : i32 | |
| %7 = vm.select.i32 %sgte, %6, %c1 : i32 | |
| %8 = vm.add.i32 %7, %c1 : i32 | |
| %9 = vm.sub.i32 %8, %c1 : i32 | |
| %10 = vm.div.i32.u %9, %c1 : i32 | |
| %11 = vm.mul.i32 %10, %9 : i32 | |
| %12 = vm.sub.i32 %7, %11 : i32 | |
| %sgte_5 = vm.cmp.gte.i32.s %12, %c1 : i32 | |
| %13 = vm.select.i32 %sgte_5, %12, %c1 : i32 | |
| %14 = vm.add.i32 %13, %c1 : i32 | |
| %15 = vm.sub.i32 %14, %c1 : i32 | |
| %16 = vm.div.i32.u %15, %c1 : i32 | |
| %zero_6 = vm.const.i32.zero : i32 | |
| vm.call.variadic @hal.command_buffer.push_constants(%ref_4, %_executable_layout_0, %zero_6, [%arg1, %arg2]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32...) | |
| %zero_7 = vm.const.i32.zero : i32 | |
| %zero_8 = vm.const.i32.zero : i32 | |
| %c1_9 = vm.const.i32 1 : i32 | |
| vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_4, %_executable_layout_0, %zero_7, [%zero_8, %c1_9], [%ref_2, %ref_1], [%zero, %zero], [%1, %1]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...) | |
| %zero_10 = vm.const.i32.zero : i32 | |
| vm.call @hal.command_buffer.dispatch(%ref_4, %_executable_simple_mul_ex_dispatch_0, %zero_10, %4, %10, %16) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> () | |
| %c20 = vm.const.i32 20 : i32 | |
| %c5 = vm.const.i32 5 : i32 | |
| %c8 = vm.const.i32 8 : i32 | |
| %c4 = vm.const.i32 4 : i32 | |
| vm.call.variadic @hal.command_buffer.execution_barrier(%ref_4, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...) | |
| vm.call @hal.command_buffer.end(%ref_4) : (!vm.ref<!hal.command_buffer>) -> () | |
| vm.call @hal.ex.submit_and_wait(%ref, %ref_4) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> () | |
| vm.return %ref_1, %arg1, %arg2 : !vm.ref<!hal.buffer>, i32, i32 | |
| } | |
| vm.export @simple_mul as("simple_mul$raw") | |
| vm.func @simple_mul$sync(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view> attributes {iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %ref = vm.call @hal.buffer_view.buffer(%arg0) : (!vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> | |
| %zero = vm.const.i32.zero : i32 | |
| %0 = vm.call @hal.buffer_view.dim(%arg0, %zero) : (!vm.ref<!hal.buffer_view>, i32) -> i32 | |
| %c1 = vm.const.i32 1 : i32 | |
| %1 = vm.call @hal.buffer_view.dim(%arg0, %c1) : (!vm.ref<!hal.buffer_view>, i32) -> i32 | |
| %ref_0:3 = vm.call @simple_mul(%ref, %0, %1) : (!vm.ref<!hal.buffer>, i32, i32) -> (!vm.ref<!hal.buffer>, i32, i32) | |
| %c50331680 = vm.const.i32 50331680 : i32 | |
| %ref_1 = vm.call.variadic @hal.buffer_view.create(%ref_0#0, [%ref_0#1, %ref_0#2], %c50331680) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view> | |
| vm.return %ref_1 : !vm.ref<!hal.buffer_view> | |
| } | |
| vm.export @simple_mul$sync as("simple_mul") | |
| vm.import @hal.ex.shared_device() -> !vm.ref<!hal.device> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.ex.defer_release(%operand : !vm.ref<?>) attributes {sym_visibility = "private"} | |
| vm.import @hal.ex.submit_and_wait(%device : !vm.ref<!hal.device>, %command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"} | |
| vm.import @hal.allocator.compute_size(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.allocator.compute_offset(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.allocator.compute_range(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.allocator.allocate(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.allocator.allocate.const(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_type : i32, %value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.allocator(%buffer : !vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator> attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.subspan(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.fill(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.read_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!iree.mutable_byte_buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.write_data(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %source_buffer : !vm.ref<!iree.byte_buffer>, %source_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.copy_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.load(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.store(%value : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %shape : i32..., %element_type : i32) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.subview(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.buffer(%buffer_view : !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.byte_length(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.compute_offset(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.compute_range(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.rank(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dim(%buffer_view : !vm.ref<!hal.buffer_view>, %index : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.1(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.2(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.3(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.4(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !vm.ref<!hal.command_buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.begin(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.end(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.fill_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.copy_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.push_constants(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %offset : i32, %values : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.push_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.bind_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %descriptor_set : !vm.ref<!hal.descriptor_set>, %dynamic_offsets : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.dispatch.indirect(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroups_buffer : !vm.ref<!hal.buffer>, %workgroups_offset : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.descriptor_set.create(%device : !vm.ref<!hal.device>, %set_layout : !vm.ref<!hal.descriptor_set_layout>, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) -> !vm.ref<!hal.descriptor_set> attributes {sym_visibility = "private"} | |
| vm.import @hal.descriptor_set_layout.create(%device : !vm.ref<!hal.device>, %usage_type : i32, %bindings : tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_cache.create(%device : !vm.ref<!hal.device>, %identifier : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_cache.select_format(%executable_cache : !vm.ref<!hal.executable_cache>, %available_formats : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_cache.prepare(%executable_cache : !vm.ref<!hal.executable_cache>, %executable_layout : !vm.ref<!hal.executable_layout>, %caching_mode : i32, %executable_data : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_layout.create(%device : !vm.ref<!hal.device>, %set_layouts : !vm.ref<!hal.descriptor_set_layout>..., %push_constants : i32) -> !vm.ref<!hal.executable_layout> attributes {nosideeffects, sym_visibility = "private"} | |
| } | |
| *** IR Dump Before CSE *** | |
| module { | |
| vm.module @module { | |
| vm.global.ref @_executable_simple_mul_ex_dispatch_0 mutable : !vm.ref<!hal.executable> | |
| vm.global.ref @_descriptor_set_layout_0 mutable : !vm.ref<!hal.descriptor_set_layout> | |
| vm.func @_descriptor_set_layout_0_initializer() -> !vm.ref<!hal.descriptor_set_layout> attributes {sym_visibility = "private"} { | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %c1 = vm.const.i32 1 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %c7 = vm.const.i32 7 : i32 | |
| %c1_0 = vm.const.i32 1 : i32 | |
| %c1_1 = vm.const.i32 1 : i32 | |
| %c7_2 = vm.const.i32 7 : i32 | |
| %c6 = vm.const.i32 6 : i32 | |
| %ref_3 = vm.call.variadic @hal.descriptor_set_layout.create(%ref, %c1, [%zero, %c7, %c1_0, %c1_1, %c7_2, %c6]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> | |
| vm.return %ref_3 : !vm.ref<!hal.descriptor_set_layout> | |
| } | |
| vm.global.ref @_executable_layout_0 mutable : !vm.ref<!hal.executable_layout> | |
| vm.func @_executable_layout_0_initializer() -> !vm.ref<!hal.executable_layout> attributes {sym_visibility = "private"} { | |
| %_descriptor_set_layout_0 = vm.global.load.ref @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout> | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %c2 = vm.const.i32 2 : i32 | |
| %ref_0 = vm.call.variadic @hal.executable_layout.create(%ref, [%_descriptor_set_layout_0], %c2) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout> | |
| vm.return %ref_0 : !vm.ref<!hal.executable_layout> | |
| } | |
| vm.global.ref @_executable_cache mutable : !vm.ref<!hal.executable_cache> | |
| vm.rodata @_utf8_default_7FD5254DFCA3A5D0 dense<[100, 101, 102, 97, 117, 108, 116]> : vector<7xi8> | |
| vm.rodata @_simple_mul_ex_dispatch_0_binary_vmla dense<"0x10000000564D4C41000006000800040006000000040000009C03000020000000424D4F44180020001C001800140010000C000000000000000800040018000000D4020000600000005400000048000000280000001400000004000000060000006D6F64756C6500000300000098020000740200005402000006000000FC010000AC010000680100001C010000D40000009C000000010000005800000001000000140000000100000000000000A00000000600040030FEFFFF0800000014000000C2FDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F30000000006CFEFFFF0800000014000000FEFDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000A8FEFFFF08000000180000003AFEFFFF040000000200000001000000010000000C000000766D6C612E6162732E66333200000000DCFEFFFF08000000240000006EFEFFFF0400000005000000010000000000000001000000000000000000000010000000766D6C612E6275666665722E636F70790000000020FFFFFF08000000280000002CFFFFFF080000000C00000001000000010000000300000001000000000000000000000010000000766D6C612E6275666665722E766965770000000068FFFFFF080000002000000074FFFFFF080000000C0000000100000001000000010000000000000011000000766D6C612E6275666665722E616C6C6F63000000A8FFFFFF0800000028000000B4FFFFFF080000000C00000001000000010000000300000002000000000000000000000016000000766D6C612E696E746572666163652E62696E64696E670000F4FFFFFF100000002C00000008000C000800040008000000080000000C000000010000000000000002000000020000000000000014000000766D6C612E696E746572666163652E636F6E737400000000CAFFFFFF040000000F00000021766D6C612E696E7465726661636500E6FFFFFF040000000C00000021766D6C612E627566666572000006000800040006000000040000000300000069333200A0000000090400000000000801000901000000020052000000800200008002000100030052000000800200008001000100040052010000800300008001000100010001802404000000000024000003000000520200008001000000010002805203000080030001C001000000010001805205000080020001C0028000005201000080030000C001000200010000805204000080050002C0010000C0010000000000540000"> : vector<952xi8> | |
| vm.func @_executable_cache_initializer() -> !vm.ref<!hal.executable_cache> attributes {sym_visibility = "private"} { | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %_utf8_default_7FD5254DFCA3A5D0 = vm.const.ref.rodata @_utf8_default_7FD5254DFCA3A5D0 : !vm.ref<!iree.byte_buffer> | |
| %ref_0 = vm.call @hal.executable_cache.create(%ref, %_utf8_default_7FD5254DFCA3A5D0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> | |
| %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout> | |
| %c1447906369 = vm.const.i32 1447906369 : i32 | |
| %0 = vm.call.variadic @hal.executable_cache.select_format(%ref_0, [%c1447906369]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32 | |
| %_simple_mul_ex_dispatch_0_binary_vmla = vm.const.ref.rodata @_simple_mul_ex_dispatch_0_binary_vmla : !vm.ref<!iree.byte_buffer> | |
| %null = vm.const.ref.zero : !vm.ref<!iree.byte_buffer> | |
| %ref_1 = vm.switch.ref %0[%_simple_mul_ex_dispatch_0_binary_vmla] else %null : !vm.ref<!iree.byte_buffer> | |
| %c7 = vm.const.i32 7 : i32 | |
| %ref_2 = vm.call @hal.executable_cache.prepare(%ref_0, %_executable_layout_0, %c7, %ref_1) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> | |
| vm.global.store.ref %ref_2, @_executable_simple_mul_ex_dispatch_0 : !vm.ref<!hal.executable> | |
| vm.return %ref_0 : !vm.ref<!hal.executable_cache> | |
| } | |
| vm.func @simple_mul(%arg0: !vm.ref<!hal.buffer>, %arg1: i32, %arg2: i32) -> (!vm.ref<!hal.buffer>, i32, i32) { | |
| %c1 = vm.const.i32 1 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %0 = vm.mul.i32 %arg1, %arg2 : i32 | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %ref_0 = vm.call @hal.device.allocator(%ref) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> | |
| %c50331680 = vm.const.i32 50331680 : i32 | |
| %1 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [%arg1, %arg2], %c50331680) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32 | |
| %c50 = vm.const.i32 50 : i32 | |
| %c15 = vm.const.i32 15 : i32 | |
| %ref_1 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %1) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer> | |
| vm.call @hal.ex.defer_release(%ref_1) : (!vm.ref<!hal.buffer>) -> () | |
| %c48 = vm.const.i32 48 : i32 | |
| %c10 = vm.const.i32 10 : i32 | |
| %ref_2 = vm.call @hal.allocator.allocate(%ref_0, %c48, %c10, %1) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer> | |
| vm.call @hal.ex.defer_release(%ref_2) : (!vm.ref<!hal.buffer>) -> () | |
| %c1_3 = vm.const.i32 1 : i32 | |
| %c3 = vm.const.i32 3 : i32 | |
| %ref_4 = vm.call @hal.command_buffer.create(%ref, %c1_3, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer> | |
| vm.call @hal.command_buffer.begin(%ref_4) : (!vm.ref<!hal.command_buffer>) -> () | |
| %_executable_simple_mul_ex_dispatch_0 = vm.global.load.ref @_executable_simple_mul_ex_dispatch_0 : !vm.ref<!hal.executable> | |
| %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout> | |
| %2 = vm.add.i32 %0, %c1 : i32 | |
| %3 = vm.sub.i32 %2, %c1 : i32 | |
| %4 = vm.div.i32.u %3, %c1 : i32 | |
| %5 = vm.mul.i32 %4, %3 : i32 | |
| %6 = vm.sub.i32 %0, %5 : i32 | |
| %sgte = vm.cmp.gte.i32.s %6, %c1 : i32 | |
| %7 = vm.select.i32 %sgte, %6, %c1 : i32 | |
| %8 = vm.add.i32 %7, %c1 : i32 | |
| %9 = vm.sub.i32 %8, %c1 : i32 | |
| %10 = vm.div.i32.u %9, %c1 : i32 | |
| %11 = vm.mul.i32 %10, %9 : i32 | |
| %12 = vm.sub.i32 %7, %11 : i32 | |
| %sgte_5 = vm.cmp.gte.i32.s %12, %c1 : i32 | |
| %13 = vm.select.i32 %sgte_5, %12, %c1 : i32 | |
| %14 = vm.add.i32 %13, %c1 : i32 | |
| %15 = vm.sub.i32 %14, %c1 : i32 | |
| %16 = vm.div.i32.u %15, %c1 : i32 | |
| %zero_6 = vm.const.i32.zero : i32 | |
| vm.call.variadic @hal.command_buffer.push_constants(%ref_4, %_executable_layout_0, %zero_6, [%arg1, %arg2]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32...) | |
| %zero_7 = vm.const.i32.zero : i32 | |
| %zero_8 = vm.const.i32.zero : i32 | |
| %c1_9 = vm.const.i32 1 : i32 | |
| vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_4, %_executable_layout_0, %zero_7, [%zero_8, %c1_9], [%ref_2, %ref_1], [%zero, %zero], [%1, %1]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...) | |
| %zero_10 = vm.const.i32.zero : i32 | |
| vm.call @hal.command_buffer.dispatch(%ref_4, %_executable_simple_mul_ex_dispatch_0, %zero_10, %4, %10, %16) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> () | |
| %c20 = vm.const.i32 20 : i32 | |
| %c5 = vm.const.i32 5 : i32 | |
| %c8 = vm.const.i32 8 : i32 | |
| %c4 = vm.const.i32 4 : i32 | |
| vm.call.variadic @hal.command_buffer.execution_barrier(%ref_4, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...) | |
| vm.call @hal.command_buffer.end(%ref_4) : (!vm.ref<!hal.command_buffer>) -> () | |
| vm.call @hal.ex.submit_and_wait(%ref, %ref_4) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> () | |
| vm.return %ref_1, %arg1, %arg2 : !vm.ref<!hal.buffer>, i32, i32 | |
| } | |
| vm.export @simple_mul as("simple_mul$raw") | |
| vm.func @simple_mul$sync(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view> attributes {iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %ref = vm.call @hal.buffer_view.buffer(%arg0) : (!vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> | |
| %zero = vm.const.i32.zero : i32 | |
| %0 = vm.call @hal.buffer_view.dim(%arg0, %zero) : (!vm.ref<!hal.buffer_view>, i32) -> i32 | |
| %c1 = vm.const.i32 1 : i32 | |
| %1 = vm.call @hal.buffer_view.dim(%arg0, %c1) : (!vm.ref<!hal.buffer_view>, i32) -> i32 | |
| %ref_0:3 = vm.call @simple_mul(%ref, %0, %1) : (!vm.ref<!hal.buffer>, i32, i32) -> (!vm.ref<!hal.buffer>, i32, i32) | |
| %c50331680 = vm.const.i32 50331680 : i32 | |
| %ref_1 = vm.call.variadic @hal.buffer_view.create(%ref_0#0, [%ref_0#1, %ref_0#2], %c50331680) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view> | |
| vm.return %ref_1 : !vm.ref<!hal.buffer_view> | |
| } | |
| vm.export @simple_mul$sync as("simple_mul") | |
| vm.import @hal.ex.shared_device() -> !vm.ref<!hal.device> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.ex.defer_release(%operand : !vm.ref<?>) attributes {sym_visibility = "private"} | |
| vm.import @hal.ex.submit_and_wait(%device : !vm.ref<!hal.device>, %command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"} | |
| vm.import @hal.allocator.compute_size(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.allocator.compute_offset(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.allocator.compute_range(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.allocator.allocate(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.allocator.allocate.const(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_type : i32, %value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.allocator(%buffer : !vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator> attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.subspan(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.fill(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.read_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!iree.mutable_byte_buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.write_data(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %source_buffer : !vm.ref<!iree.byte_buffer>, %source_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.copy_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.load(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.store(%value : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %shape : i32..., %element_type : i32) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.subview(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.buffer(%buffer_view : !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.byte_length(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.compute_offset(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.compute_range(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.rank(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dim(%buffer_view : !vm.ref<!hal.buffer_view>, %index : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.1(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.2(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.3(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.4(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !vm.ref<!hal.command_buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.begin(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.end(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.fill_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.copy_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.push_constants(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %offset : i32, %values : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.push_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.bind_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %descriptor_set : !vm.ref<!hal.descriptor_set>, %dynamic_offsets : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.dispatch.indirect(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroups_buffer : !vm.ref<!hal.buffer>, %workgroups_offset : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.descriptor_set.create(%device : !vm.ref<!hal.device>, %set_layout : !vm.ref<!hal.descriptor_set_layout>, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) -> !vm.ref<!hal.descriptor_set> attributes {sym_visibility = "private"} | |
| vm.import @hal.descriptor_set_layout.create(%device : !vm.ref<!hal.device>, %usage_type : i32, %bindings : tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_cache.create(%device : !vm.ref<!hal.device>, %identifier : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_cache.select_format(%executable_cache : !vm.ref<!hal.executable_cache>, %available_formats : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_cache.prepare(%executable_cache : !vm.ref<!hal.executable_cache>, %executable_layout : !vm.ref<!hal.executable_layout>, %caching_mode : i32, %executable_data : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_layout.create(%device : !vm.ref<!hal.device>, %set_layouts : !vm.ref<!hal.descriptor_set_layout>..., %push_constants : i32) -> !vm.ref<!hal.executable_layout> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.func @__init() { | |
| %ref = vm.call @_descriptor_set_layout_0_initializer() : () -> !vm.ref<!hal.descriptor_set_layout> | |
| vm.global.store.ref %ref, @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout> | |
| %ref_0 = vm.call @_executable_layout_0_initializer() : () -> !vm.ref<!hal.executable_layout> | |
| vm.global.store.ref %ref_0, @_executable_layout_0 : !vm.ref<!hal.executable_layout> | |
| %ref_1 = vm.call @_executable_cache_initializer() : () -> !vm.ref<!hal.executable_cache> | |
| vm.global.store.ref %ref_1, @_executable_cache : !vm.ref<!hal.executable_cache> | |
| vm.return | |
| } | |
| vm.export @__init | |
| } | |
| } | |
| *** IR Dump Before SymbolDCE *** | |
| module { | |
| vm.module @module { | |
| vm.global.ref @_executable_simple_mul_ex_dispatch_0 mutable : !vm.ref<!hal.executable> | |
| vm.global.ref @_descriptor_set_layout_0 mutable : !vm.ref<!hal.descriptor_set_layout> | |
| vm.func @_descriptor_set_layout_0_initializer() -> !vm.ref<!hal.descriptor_set_layout> attributes {sym_visibility = "private"} { | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %c1 = vm.const.i32 1 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %c7 = vm.const.i32 7 : i32 | |
| %c6 = vm.const.i32 6 : i32 | |
| %ref_0 = vm.call.variadic @hal.descriptor_set_layout.create(%ref, %c1, [%zero, %c7, %c1, %c1, %c7, %c6]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> | |
| vm.return %ref_0 : !vm.ref<!hal.descriptor_set_layout> | |
| } | |
| vm.global.ref @_executable_layout_0 mutable : !vm.ref<!hal.executable_layout> | |
| vm.func @_executable_layout_0_initializer() -> !vm.ref<!hal.executable_layout> attributes {sym_visibility = "private"} { | |
| %_descriptor_set_layout_0 = vm.global.load.ref @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout> | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %c2 = vm.const.i32 2 : i32 | |
| %ref_0 = vm.call.variadic @hal.executable_layout.create(%ref, [%_descriptor_set_layout_0], %c2) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout> | |
| vm.return %ref_0 : !vm.ref<!hal.executable_layout> | |
| } | |
| vm.global.ref @_executable_cache mutable : !vm.ref<!hal.executable_cache> | |
| vm.rodata @_utf8_default_7FD5254DFCA3A5D0 dense<[100, 101, 102, 97, 117, 108, 116]> : vector<7xi8> | |
| vm.rodata @_simple_mul_ex_dispatch_0_binary_vmla dense<"0x10000000564D4C41000006000800040006000000040000009C03000020000000424D4F44180020001C001800140010000C000000000000000800040018000000D4020000600000005400000048000000280000001400000004000000060000006D6F64756C6500000300000098020000740200005402000006000000FC010000AC010000680100001C010000D40000009C000000010000005800000001000000140000000100000000000000A00000000600040030FEFFFF0800000014000000C2FDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F30000000006CFEFFFF0800000014000000FEFDFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000A8FEFFFF08000000180000003AFEFFFF040000000200000001000000010000000C000000766D6C612E6162732E66333200000000DCFEFFFF08000000240000006EFEFFFF0400000005000000010000000000000001000000000000000000000010000000766D6C612E6275666665722E636F70790000000020FFFFFF08000000280000002CFFFFFF080000000C00000001000000010000000300000001000000000000000000000010000000766D6C612E6275666665722E766965770000000068FFFFFF080000002000000074FFFFFF080000000C0000000100000001000000010000000000000011000000766D6C612E6275666665722E616C6C6F63000000A8FFFFFF0800000028000000B4FFFFFF080000000C00000001000000010000000300000002000000000000000000000016000000766D6C612E696E746572666163652E62696E64696E670000F4FFFFFF100000002C00000008000C000800040008000000080000000C000000010000000000000002000000020000000000000014000000766D6C612E696E746572666163652E636F6E737400000000CAFFFFFF040000000F00000021766D6C612E696E7465726661636500E6FFFFFF040000000C00000021766D6C612E627566666572000006000800040006000000040000000300000069333200A0000000090400000000000801000901000000020052000000800200008002000100030052000000800200008001000100040052010000800300008001000100010001802404000000000024000003000000520200008001000000010002805203000080030001C001000000010001805205000080020001C0028000005201000080030000C001000200010000805204000080050002C0010000C0010000000000540000"> : vector<952xi8> | |
| vm.func @_executable_cache_initializer() -> !vm.ref<!hal.executable_cache> attributes {sym_visibility = "private"} { | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %_utf8_default_7FD5254DFCA3A5D0 = vm.const.ref.rodata @_utf8_default_7FD5254DFCA3A5D0 : !vm.ref<!iree.byte_buffer> | |
| %ref_0 = vm.call @hal.executable_cache.create(%ref, %_utf8_default_7FD5254DFCA3A5D0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> | |
| %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout> | |
| %c1447906369 = vm.const.i32 1447906369 : i32 | |
| %0 = vm.call.variadic @hal.executable_cache.select_format(%ref_0, [%c1447906369]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32 | |
| %_simple_mul_ex_dispatch_0_binary_vmla = vm.const.ref.rodata @_simple_mul_ex_dispatch_0_binary_vmla : !vm.ref<!iree.byte_buffer> | |
| %null = vm.const.ref.zero : !vm.ref<!iree.byte_buffer> | |
| %ref_1 = vm.switch.ref %0[%_simple_mul_ex_dispatch_0_binary_vmla] else %null : !vm.ref<!iree.byte_buffer> | |
| %c7 = vm.const.i32 7 : i32 | |
| %ref_2 = vm.call @hal.executable_cache.prepare(%ref_0, %_executable_layout_0, %c7, %ref_1) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> | |
| vm.global.store.ref %ref_2, @_executable_simple_mul_ex_dispatch_0 : !vm.ref<!hal.executable> | |
| vm.return %ref_0 : !vm.ref<!hal.executable_cache> | |
| } | |
| vm.func @simple_mul(%arg0: !vm.ref<!hal.buffer>, %arg1: i32, %arg2: i32) -> (!vm.ref<!hal.buffer>, i32, i32) { | |
| %c1 = vm.const.i32 1 : i32 | |
| %zero = vm.const.i32.zero : i32 | |
| %0 = vm.mul.i32 %arg1, %arg2 : i32 | |
| %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device> | |
| %ref_0 = vm.call @hal.device.allocator(%ref) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> | |
| %c50331680 = vm.const.i32 50331680 : i32 | |
| %1 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [%arg1, %arg2], %c50331680) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32 | |
| %c50 = vm.const.i32 50 : i32 | |
| %c15 = vm.const.i32 15 : i32 | |
| %ref_1 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %1) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer> | |
| vm.call @hal.ex.defer_release(%ref_1) : (!vm.ref<!hal.buffer>) -> () | |
| %c48 = vm.const.i32 48 : i32 | |
| %c10 = vm.const.i32 10 : i32 | |
| %ref_2 = vm.call @hal.allocator.allocate(%ref_0, %c48, %c10, %1) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer> | |
| vm.call @hal.ex.defer_release(%ref_2) : (!vm.ref<!hal.buffer>) -> () | |
| %c3 = vm.const.i32 3 : i32 | |
| %ref_3 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer> | |
| vm.call @hal.command_buffer.begin(%ref_3) : (!vm.ref<!hal.command_buffer>) -> () | |
| %_executable_simple_mul_ex_dispatch_0 = vm.global.load.ref @_executable_simple_mul_ex_dispatch_0 : !vm.ref<!hal.executable> | |
| %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout> | |
| %2 = vm.add.i32 %0, %c1 : i32 | |
| %3 = vm.sub.i32 %2, %c1 : i32 | |
| %4 = vm.div.i32.u %3, %c1 : i32 | |
| %5 = vm.mul.i32 %4, %3 : i32 | |
| %6 = vm.sub.i32 %0, %5 : i32 | |
| %sgte = vm.cmp.gte.i32.s %6, %c1 : i32 | |
| %7 = vm.select.i32 %sgte, %6, %c1 : i32 | |
| %8 = vm.add.i32 %7, %c1 : i32 | |
| %9 = vm.sub.i32 %8, %c1 : i32 | |
| %10 = vm.div.i32.u %9, %c1 : i32 | |
| %11 = vm.mul.i32 %10, %9 : i32 | |
| %12 = vm.sub.i32 %7, %11 : i32 | |
| %sgte_4 = vm.cmp.gte.i32.s %12, %c1 : i32 | |
| %13 = vm.select.i32 %sgte_4, %12, %c1 : i32 | |
| %14 = vm.add.i32 %13, %c1 : i32 | |
| %15 = vm.sub.i32 %14, %c1 : i32 | |
| %16 = vm.div.i32.u %15, %c1 : i32 | |
| vm.call.variadic @hal.command_buffer.push_constants(%ref_3, %_executable_layout_0, %zero, [%arg1, %arg2]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32...) | |
| vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_3, %_executable_layout_0, %zero, [%zero, %c1], [%ref_2, %ref_1], [%zero, %zero], [%1, %1]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...) | |
| vm.call @hal.command_buffer.dispatch(%ref_3, %_executable_simple_mul_ex_dispatch_0, %zero, %4, %10, %16) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> () | |
| %c20 = vm.const.i32 20 : i32 | |
| %c5 = vm.const.i32 5 : i32 | |
| %c8 = vm.const.i32 8 : i32 | |
| %c4 = vm.const.i32 4 : i32 | |
| vm.call.variadic @hal.command_buffer.execution_barrier(%ref_3, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...) | |
| vm.call @hal.command_buffer.end(%ref_3) : (!vm.ref<!hal.command_buffer>) -> () | |
| vm.call @hal.ex.submit_and_wait(%ref, %ref_3) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> () | |
| vm.return %ref_1, %arg1, %arg2 : !vm.ref<!hal.buffer>, i32, i32 | |
| } | |
| vm.export @simple_mul as("simple_mul$raw") | |
| vm.func @simple_mul$sync(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view> attributes {iree.reflection = {f = "I10!B7!d-1d-1R10!B7!d-1d-1", fv = "1"}} { | |
| %ref = vm.call @hal.buffer_view.buffer(%arg0) : (!vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> | |
| %zero = vm.const.i32.zero : i32 | |
| %0 = vm.call @hal.buffer_view.dim(%arg0, %zero) : (!vm.ref<!hal.buffer_view>, i32) -> i32 | |
| %c1 = vm.const.i32 1 : i32 | |
| %1 = vm.call @hal.buffer_view.dim(%arg0, %c1) : (!vm.ref<!hal.buffer_view>, i32) -> i32 | |
| %ref_0:3 = vm.call @simple_mul(%ref, %0, %1) : (!vm.ref<!hal.buffer>, i32, i32) -> (!vm.ref<!hal.buffer>, i32, i32) | |
| %c50331680 = vm.const.i32 50331680 : i32 | |
| %ref_1 = vm.call.variadic @hal.buffer_view.create(%ref_0#0, [%ref_0#1, %ref_0#2], %c50331680) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view> | |
| vm.return %ref_1 : !vm.ref<!hal.buffer_view> | |
| } | |
| vm.export @simple_mul$sync as("simple_mul") | |
| vm.import @hal.ex.shared_device() -> !vm.ref<!hal.device> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.ex.defer_release(%operand : !vm.ref<?>) attributes {sym_visibility = "private"} | |
| vm.import @hal.ex.submit_and_wait(%device : !vm.ref<!hal.device>, %command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"} | |
| vm.import @hal.allocator.compute_size(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.allocator.compute_offset(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.allocator.compute_range(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.allocator.allocate(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.allocator.allocate.const(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_type : i32, %value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.allocator(%buffer : !vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator> attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.subspan(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.fill(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.read_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!iree.mutable_byte_buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.write_data(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %source_buffer : !vm.ref<!iree.byte_buffer>, %source_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.copy_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.load(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer.store(%value : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %shape : i32..., %element_type : i32) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.subview(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.buffer(%buffer_view : !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.byte_length(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.compute_offset(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.compute_range(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.rank(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dim(%buffer_view : !vm.ref<!hal.buffer_view>, %index : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.1(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.2(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.3(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.buffer_view.dims.4(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !vm.ref<!hal.command_buffer> attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.begin(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.end(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.fill_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.copy_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.push_constants(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %offset : i32, %values : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.push_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.bind_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %descriptor_set : !vm.ref<!hal.descriptor_set>, %dynamic_offsets : i32...) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.command_buffer.dispatch.indirect(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroups_buffer : !vm.ref<!hal.buffer>, %workgroups_offset : i32) attributes {sym_visibility = "private"} | |
| vm.import @hal.descriptor_set.create(%device : !vm.ref<!hal.device>, %set_layout : !vm.ref<!hal.descriptor_set_layout>, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) -> !vm.ref<!hal.descriptor_set> attributes {sym_visibility = "private"} | |
| vm.import @hal.descriptor_set_layout.create(%device : !vm.ref<!hal.device>, %usage_type : i32, %bindings : tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_cache.create(%device : !vm.ref<!hal.device>, %identifier : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_cache.select_format(%executable_cache : !vm.ref<!hal.executable_cache>, %available_formats : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_cache.prepare(%executable_cache : !vm.ref<!hal.executable_cache>, %executable_layout : !vm.ref<!hal.executable_layout>, %caching_mode : i32, %executable_data : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.import @hal.executable_layout.create(%device : !vm.ref<!hal.device>, %set_layouts : !vm.ref<!hal.descriptor_set_layout>..., %push_constants : i32) -> !vm.ref<!hal.executable_layout> attributes {nosideeffects, sym_visibility = "private"} | |
| vm.func @__init() { | |
| %ref = vm.call @_descriptor_set_layout_0_initializer() : () -> !vm.ref<!hal.descriptor_set_layout> | |
| vm.global.store.ref %ref, @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout> | |
| %ref_0 = vm.call @_executable_layout_0_initializer() : () -> !vm.ref<!hal.executable_layout> | |
| vm.global.store.ref %ref_0, @_executable_layout_0 : !vm.ref<!hal.executable_layout> | |
| %ref_1 = vm.call @_executable_cache_initializer() : () -> !vm.ref<!hal.executable_cache> | |
| vm.global.store.ref %ref_1, @_executable_cache : !vm.ref<!hal.executable_cache> | |
| vm.return | |
| } | |
| vm.export @__init | |
| } | |
| } | |
| WARNING: Logging before InitGoogle() is written to STDERR | |
| E0329 20:52:24.897359 58482 api_util.h:30] INVALID_ARGUMENT: while executing hal.command_buffer.push_constants | |
| [ FAILED ] VmTest.test_synchronous_dynamic_shape_invoke_function | |
| ====================================================================== | |
| ERROR: test_synchronous_dynamic_shape_invoke_function (__main__.VmTest) | |
| test_synchronous_dynamic_shape_invoke_function (__main__.VmTest) | |
| ---------------------------------------------------------------------- | |
| Traceback (most recent call last): | |
| File "/usr/local/google/_blaze_laurenzo/76d22ed88fb985a166e4e73d755722b5/execroot/google3/blaze-out/k8-dbg/bin/third_party/iree/bindings/python/pyiree/rt/vm_test.runfiles/google3/third_party/iree/bindings/python/pyiree/rt/vm_test.py", line 120, in test_synchronous_dynamic_shape_invoke_function | |
| context.invoke(f, inputs, allocated_results) | |
| ValueError: Error invoking function: 3 | |
| ---------------------------------------------------------------------- | |
| Ran 1 test in 7.319s | |
| FAILED (errors=1) | |
| DRIVER_NAMES = ['llvm', 'vmla', 'vulkan'] | |
| INVOKING: <FunctionAbi (Buffer<float32[?x?]>) -> (Buffer<float32[?x?]>)> | |
| INPUTS: <VmVariantList(1): [HalBufferView(2x2:0x3000020)]> | |
| ALLOCATED RESULTS: <VmVariantList(0): []> | |
| --- INVOKE: |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment