Created
January 21, 2020 21:54
-
-
Save benvanik/1815ca5b6ab77434245b53d211f17f9b to your computer and use it in GitHub Desktop.
IREE HLO -> flatbuffer end-to-end example dump
This file contains 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 After Canonicalizer *** | |
| module @hal_usage { | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| } | |
| *** IR Dump After xla_hlo::(anonymous namespace)::LegalizeControlFlow *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::(anonymous namespace)::FlattenTuplesInCFGPass *** | |
| module @hal_usage { | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| } | |
| *** IR Dump After InlinerPass *** | |
| module @hal_usage { | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| } | |
| *** IR Dump After Canonicalizer *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| *** IR Dump After CSE *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::LegalizeInputTypesPass *** | |
| module @hal_usage { | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::PrePartitioningConversionPass *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::IdentifyReductionRegionsPass *** | |
| module @hal_usage { | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| } | |
| *** IR Dump After CSE *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::DispatchabilityAnalysisPass *** | |
| module @hal_usage { | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::IdentifyDispatchRegionsPass *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.dispatch.region[%cst : vector<3xi32>](%arg1 = %arg0 : tensor<4xf32>) -> tensor<4xf32> { | |
| %1 = xla_hlo.add %arg1, %arg1 : tensor<4xf32> | |
| %2 = xla_hlo.sub %1, %arg1 : tensor<4xf32> | |
| %3 = xla_hlo.mul %2, %arg1 : tensor<4xf32> | |
| flow.return %3 : tensor<4xf32> | |
| } | |
| return %0 : tensor<4xf32> | |
| } | |
| *** IR Dump After CSE *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.dispatch.region[%cst : vector<3xi32>](%arg1 = %arg0 : tensor<4xf32>) -> tensor<4xf32> { | |
| %1 = xla_hlo.add %arg1, %arg1 : tensor<4xf32> | |
| %2 = xla_hlo.sub %1, %arg1 : tensor<4xf32> | |
| %3 = xla_hlo.mul %2, %arg1 : tensor<4xf32> | |
| flow.return %3 : tensor<4xf32> | |
| } | |
| return %0 : tensor<4xf32> | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::FoldCompatibleDispatchRegionsPass *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.dispatch.region[%cst : vector<3xi32>](%arg1 = %arg0 : tensor<4xf32>) -> tensor<4xf32> { | |
| %1 = xla_hlo.add %arg1, %arg1 : tensor<4xf32> | |
| %2 = xla_hlo.sub %1, %arg1 : tensor<4xf32> | |
| %3 = xla_hlo.mul %2, %arg1 : tensor<4xf32> | |
| flow.return %3 : tensor<4xf32> | |
| } | |
| return %0 : tensor<4xf32> | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::RematerializeDispatchConstantsPass *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.dispatch.region[%cst : vector<3xi32>](%arg1 = %arg0 : tensor<4xf32>) -> tensor<4xf32> { | |
| %1 = xla_hlo.add %arg1, %arg1 : tensor<4xf32> | |
| %2 = xla_hlo.sub %1, %arg1 : tensor<4xf32> | |
| %3 = xla_hlo.mul %2, %arg1 : tensor<4xf32> | |
| flow.return %3 : tensor<4xf32> | |
| } | |
| return %0 : tensor<4xf32> | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::OutlineDispatchRegionsPass *** | |
| module @hal_usage { | |
| flow.executable @hloElementwiseOps_ex_dispatch_0 { | |
| flow.dispatch.entry @hloElementwiseOps_rgn_dispatch_0 | |
| module { | |
| func @hloElementwiseOps_rgn_dispatch_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| } | |
| } | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.dispatch @hloElementwiseOps_ex_dispatch_0::@hloElementwiseOps_rgn_dispatch_0[%cst : vector<3xi32>](%arg0) : (tensor<4xf32>) -> tensor<4xf32> | |
| return %0 : tensor<4xf32> | |
| } | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::OutlineReductionRegionsPass *** | |
| module @hal_usage { | |
| flow.executable @hloElementwiseOps_ex_dispatch_0 { | |
| flow.dispatch.entry @hloElementwiseOps_rgn_dispatch_0 | |
| module { | |
| func @hloElementwiseOps_rgn_dispatch_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| } | |
| } | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.dispatch @hloElementwiseOps_ex_dispatch_0::@hloElementwiseOps_rgn_dispatch_0[%cst : vector<3xi32>](%arg0) : (tensor<4xf32>) -> tensor<4xf32> | |
| return %0 : tensor<4xf32> | |
| } | |
| } | |
| *** IR Dump After Canonicalizer *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.dispatch @hloElementwiseOps_ex_dispatch_0::@hloElementwiseOps_rgn_dispatch_0[%cst : vector<3xi32>](%arg0) : (tensor<4xf32>) -> tensor<4xf32> | |
| return %0 : tensor<4xf32> | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::PostPartitioningConversionPass *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.dispatch @hloElementwiseOps_ex_dispatch_0::@hloElementwiseOps_rgn_dispatch_0[%cst : vector<3xi32>](%arg0) : (tensor<4xf32>) -> tensor<4xf32> | |
| return %0 : tensor<4xf32> | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::AssignExecutableWorkloadsPass *** | |
| module @hal_usage { | |
| flow.executable @hloElementwiseOps_ex_dispatch_0 { | |
| flow.dispatch.entry @hloElementwiseOps_rgn_dispatch_0 attributes {workload = dense<[4, 1, 1]> : vector<3xi32>} | |
| module { | |
| func @hloElementwiseOps_rgn_dispatch_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| } | |
| } | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.dispatch @hloElementwiseOps_ex_dispatch_0::@hloElementwiseOps_rgn_dispatch_0[%cst : vector<3xi32>](%arg0) : (tensor<4xf32>) -> tensor<4xf32> | |
| return %0 : tensor<4xf32> | |
| } | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::FormStreamsPass *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.ex.stream.fragment(%arg1 = %cst : vector<3xi32>, %arg2 = %arg0 : tensor<4xf32>) -> tensor<4xf32> { | |
| %1 = flow.dispatch @hloElementwiseOps_ex_dispatch_0::@hloElementwiseOps_rgn_dispatch_0[%arg1 : vector<3xi32>](%arg2) : (tensor<4xf32>) -> tensor<4xf32> | |
| flow.return %1 : tensor<4xf32> | |
| } | |
| return %0 : tensor<4xf32> | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::MaterializeExportedReflectionPass *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32> {iree.reflection = {f_partial = "I6!B3!d4"}}) -> (tensor<4xf32> {iree.reflection = {f_partial = "R6!B3!d4"}}) attributes {iree.module.export} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.ex.stream.fragment(%arg1 = %cst : vector<3xi32>, %arg2 = %arg0 : tensor<4xf32>) -> tensor<4xf32> { | |
| %1 = flow.dispatch @hloElementwiseOps_ex_dispatch_0::@hloElementwiseOps_rgn_dispatch_0[%arg1 : vector<3xi32>](%arg2) : (tensor<4xf32>) -> tensor<4xf32> | |
| flow.return %1 : tensor<4xf32> | |
| } | |
| return %0 : tensor<4xf32> | |
| } | |
| *** IR Dump After iree_compiler::IREE::Flow::MergeExportedReflectionPass *** | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export, iree.reflection = {f = "I6!B3!d4R6!B3!d4", fv = "1"}} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.ex.stream.fragment(%arg1 = %cst : vector<3xi32>, %arg2 = %arg0 : tensor<4xf32>) -> tensor<4xf32> { | |
| %1 = flow.dispatch @hloElementwiseOps_ex_dispatch_0::@hloElementwiseOps_rgn_dispatch_0[%arg1 : vector<3xi32>](%arg2) : (tensor<4xf32>) -> tensor<4xf32> | |
| flow.return %1 : tensor<4xf32> | |
| } | |
| return %0 : tensor<4xf32> | |
| } | |
| *** IR Dump After Canonicalizer *** | |
| module @hal_usage { | |
| flow.executable @hloElementwiseOps_ex_dispatch_0 { | |
| flow.dispatch.entry @hloElementwiseOps_rgn_dispatch_0 attributes {workload = dense<[4, 1, 1]> : vector<3xi32>} | |
| module { | |
| func @hloElementwiseOps_rgn_dispatch_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { | |
| %0 = xla_hlo.add %arg0, %arg0 : tensor<4xf32> | |
| %1 = xla_hlo.sub %0, %arg0 : tensor<4xf32> | |
| %2 = xla_hlo.mul %1, %arg0 : tensor<4xf32> | |
| return %2 : tensor<4xf32> | |
| } | |
| } | |
| } | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export, iree.reflection = {f = "I6!B3!d4R6!B3!d4", fv = "1"}} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.ex.stream.fragment(%arg1 = %cst : vector<3xi32>, %arg2 = %arg0 : tensor<4xf32>) -> tensor<4xf32> { | |
| %1 = flow.dispatch @hloElementwiseOps_ex_dispatch_0::@hloElementwiseOps_rgn_dispatch_0[%arg1 : vector<3xi32>](%arg2) : (tensor<4xf32>) -> tensor<4xf32> | |
| flow.return %1 : tensor<4xf32> | |
| } | |
| return %0 : tensor<4xf32> | |
| } | |
| } | |
| *** IR Dump After iree_compiler::IREE::HAL::TranslateExecutablesPass *** | |
| module @hal_usage { | |
| hal.executable @hloElementwiseOps_ex_dispatch_0 { | |
| hal.executable.entry_point @hloElementwiseOps_rgn_dispatch_0 attributes {ordinal = 0 : i32, workgroup_size = dense<[32, 1, 1]> : vector<3xi32>} | |
| hal.executable.binary attributes {data = dense<[16, 0, 0, 0, 69, 77, 79, 68, 8, 0, 14, 0, 8, 0, 4, 0, 8, 0, 0, 0, 20, 0, 0, 0, 52, 0, 0, 0, 0, 0, 10, 0, 16, 0, 12, 0, 8, 0, 4, 0, 10, 0, 0, 0, 12, 0, 0, 0, 16, 0, 0, 0, 16, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 28, 0, 0, 0, 6, 0, 0, 0, 109, 111, 100, 117, 108, 101, 0, 0, 12, 0, 16, 0, 12, 0, 8, 0, 0, 0, 4, 0, 12, 0, 0, 0, -52, 0, 0, 0, 48, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 126, -1, -1, -1, 4, 0, 0, 0, 2, 0, 0, 0, 68, 0, 0, 0, 4, 0, 0, 0, -52, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -112, -1, -1, -1, 8, 0, 0, 0, 12, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, -20, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -66, -1, -1, -1, 32, 0, 0, 0, 8, 0, 12, 0, 11, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 1, -52, -1, -1, -1, 8, 0, 0, 0, 20, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, 8, 0, 10, 0, 9, 0, 4, 0, 8, 0, 0, 0, 12, 0, 0, 0, 0, 1, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 32, 0, 0, 0, 8, 0, 12, 0, 8, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 7, 0, 0, 0, 112, 0, 0, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 2, 0, 113, 0, 0, 0, 0, 2, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 3, 0, 115, 2, 0, 0, 0, 3, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 4, 0, 119, 3, 0, 0, 0, 4, 0, 0, 3, 1, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 5, 0, 0, 3, 1, 1, 0, 0, 0, 1, 4, 0, 0, 0, 0, 0, 0, 0, 6, 0, 54, 4, 0, 5, 0, 1, 0, 5, 0, 6, 0, 4, 0]> : vector<436xi8>, format = 1230128453 : i32} { | |
| module { | |
| func @hloElementwiseOps_rgn_dispatch_0(%arg0: memref<4xf32>, %arg1: memref<4xf32>) attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : vector<3xi32>, iree.executable.workload = dense<[4, 1, 1]> : vector<3xi32>, iree.ordinal = 0 : i32} { | |
| %0 = "iree_ll_interp.alloc_heap"() : () -> memref<4xf32> | |
| "iree_ll_interp.add_f"(%arg0, %arg0, %0) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> () | |
| %1 = "iree_ll_interp.alloc_heap"() : () -> memref<4xf32> | |
| "iree_ll_interp.sub_f"(%0, %arg0, %1) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> () | |
| %2 = "iree_ll_interp.alloc_heap"() : () -> memref<4xf32> | |
| "iree_ll_interp.mul_f"(%1, %arg0, %2) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> () | |
| %3 = "iree_ll_interp.constant"() {value = dense<0> : tensor<1xi64>} : () -> memref<1xi64> | |
| %4 = "iree_ll_interp.constant"() {value = dense<4> : tensor<1xi64>} : () -> memref<1xi64> | |
| "iree_ll_interp.dynamic_copy"(%2, %3, %arg1, %3, %4) : (memref<4xf32>, memref<1xi64>, memref<4xf32>, memref<1xi64>, memref<1xi64>) -> () | |
| iree.return | |
| } | |
| } | |
| } | |
| hal.executable.binary attributes {data = dense<[8, 0, 0, 0, 83, 80, 86, 69, -92, -1, -1, -1, 104, 4, 0, 0, 112, 0, 0, 0, 4, 0, 0, 0, -18, -1, -1, -1, 4, 0, 0, 0, 1, 0, 0, 0, 12, 0, 0, 0, 0, 0, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 52, 0, 0, 0, 16, 0, 0, 0, 12, 0, 20, 0, 4, 0, 8, 0, 12, 0, 16, 0, 12, 0, 0, 0, 1, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, 12, 0, 16, 0, 0, 0, 4, 0, 8, 0, 12, 0, 12, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, -4, 0, 0, 0, 3, 2, 35, 7, 0, 0, 1, 0, 22, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 17, 0, 2, 0, 1, 0, 0, 0, 10, 0, 11, 0, 83, 80, 86, 95, 75, 72, 82, 95, 115, 116, 111, 114, 97, 103, 101, 95, 98, 117, 102, 102, 101, 114, 95, 115, 116, 111, 114, 97, 103, 101, 95, 99, 108, 97, 115, 115, 0, 0, 0, 0, 14, 0, 3, 0, 0, 0, 0, 0, 1, 0, 0, 0, 15, 0, 13, 0, 5, 0, 0, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 4, 0, 0, 0, 16, 0, 6, 0, 14, 0, 0, 0, 17, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 5, 0, 7, 0, 4, 0, 0, 0, 103, 108, 111, 98, 97, 108, 73, 110, 118, 111, 99, 97, 116, 105, 111, 110, 73, 68, 0, 0, 5, 0, 12, 0, 10, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 48, 0, 0, 5, 0, 12, 0, 11, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 49, 0, 0, 5, 0, 11, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 71, 0, 4, 0, 4, 0, 0, 0, 11, 0, 0, 0, 28, 0, 0, 0, 71, 0, 4, 0, 7, 0, 0, 0, 6, 0, 0, 0, 4, 0, 0, 0, 72, 0, 5, 0, 6, 0, 0, 0, 0, 0, 0, 0, 35, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 33, 0, 0, 0, 1, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 21, 0, 4, 0, 3, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 23, 0, 4, 0, 2, 0, 0, 0, 3, 0, 0, 0, 3, 0, 0, 0, 32, 0, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 2, 0, 0, 0, 59, 0, 4, 0, 1, 0, 0, 0, 4, 0, 0, 0, 1, 0, 0, 0, 22, 0, 3, 0, 8, 0, 0, 0, 32, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 9, 0, 0, 0, 4, 0, 0, 0, 28, 0, 4, 0, 7, 0, 0, 0, 8, 0, 0, 0, 9, 0, 0, 0, 30, 0, 3, 0, 6, 0, 0, 0, 7, 0, 0, 0, 32, 0, 4, 0, 5, 0, 0, 0, 12, 0, 0, 0, 6, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 10, 0, 0, 0, 12, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 11, 0, 0, 0, 12, 0, 0, 0, 19, 0, 2, 0, 13, 0, 0, 0, 33, 0, 3, 0, 12, 0, 0, 0, 13, 0, 0, 0, 20, 0, 2, 0, 16, 0, 0, 0, 41, 0, 3, 0, 16, 0, 0, 0, 17, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 18, 0, 0, 0, 0, 0, 0, 0, 32, 0, 4, 0, 26, 0, 0, 0, 12, 0, 0, 0, 8, 0, 0, 0, 54, 0, 5, 0, 13, 0, 0, 0, 14, 0, 0, 0, 0, 0, 0, 0, 12, 0, 0, 0, -8, 0, 2, 0, 15, 0, 0, 0, 61, 0, 4, 0, 2, 0, 0, 0, 19, 0, 0, 0, 4, 0, 0, 0, 81, 0, 5, 0, 3, 0, 0, 0, 20, 0, 0, 0, 19, 0, 0, 0, 0, 0, 0, 0, -79, 0, 5, 0, 16, 0, 0, 0, 21, 0, 0, 0, 20, 0, 0, 0, 9, 0, 0, 0, -89, 0, 5, 0, 16, 0, 0, 0, 22, 0, 0, 0, 21, 0, 0, 0, 17, 0, 0, 0, -9, 0, 3, 0, 25, 0, 0, 0, 0, 0, 0, 0, -6, 0, 4, 0, 22, 0, 0, 0, 24, 0, 0, 0, 25, 0, 0, 0, -8, 0, 2, 0, 24, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 27, 0, 0, 0, 10, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 61, 0, 4, 0, 8, 0, 0, 0, 28, 0, 0, 0, 27, 0, 0, 0, -127, 0, 5, 0, 8, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, 28, 0, 0, 0, -125, 0, 5, 0, 8, 0, 0, 0, 30, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, -123, 0, 5, 0, 8, 0, 0, 0, 31, 0, 0, 0, 30, 0, 0, 0, 28, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 32, 0, 0, 0, 11, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 62, 0, 3, 0, 32, 0, 0, 0, 31, 0, 0, 0, -7, 0, 2, 0, 25, 0, 0, 0, -8, 0, 2, 0, 25, 0, 0, 0, -3, 0, 1, 0, 56, 0, 1, 0, 1, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0]> : vector<1188xi8>, format = 1397773893 : i32} { | |
| module { | |
| spv.module "Logical" "GLSL450" { | |
| spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> | |
| spv.globalVariable @hloElementwiseOps_rgn_dispatch_0_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| spv.globalVariable @hloElementwiseOps_rgn_dispatch_0_arg_1 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| func @hloElementwiseOps_rgn_dispatch_0() { | |
| %0 = spv.constant true | |
| %1 = spv.constant 4 : i32 | |
| %2 = spv.constant 0 : i32 | |
| %3 = spv._address_of @hloElementwiseOps_rgn_dispatch_0_arg_1 : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| %4 = spv._address_of @hloElementwiseOps_rgn_dispatch_0_arg_0 : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| %5 = spv._address_of @globalInvocationID : !spv.ptr<vector<3xi32>, Input> | |
| %6 = spv.Load "Input" %5 : vector<3xi32> | |
| %7 = spv.CompositeExtract %6[0 : i32] : vector<3xi32> | |
| %8 = spv.SLessThan %7, %1 : i32 | |
| %9 = spv.LogicalAnd %8, %0 : i1 | |
| spv.selection { | |
| spv.BranchConditional %9, ^bb1, ^bb2 | |
| ^bb1: // pred: ^bb0 | |
| %10 = spv.AccessChain %4[%2, %7] : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| %11 = spv.Load "StorageBuffer" %10 : f32 | |
| %12 = spv.FAdd %11, %11 : f32 | |
| %13 = spv.FSub %12, %11 : f32 | |
| %14 = spv.FMul %13, %11 : f32 | |
| %15 = spv.AccessChain %3[%2, %7] : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| spv.Store "StorageBuffer" %15, %14 : f32 | |
| spv.Branch ^bb2 | |
| ^bb2: // 2 preds: ^bb0, ^bb1 | |
| spv._merge | |
| } | |
| spv.Return | |
| } | |
| spv.EntryPoint "GLCompute" @hloElementwiseOps_rgn_dispatch_0, @globalInvocationID | |
| spv.ExecutionMode @hloElementwiseOps_rgn_dispatch_0 "LocalSize", 32, 1, 1 | |
| } attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]} | |
| } | |
| } | |
| } | |
| func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export, iree.reflection = {f = "I6!B3!d4R6!B3!d4", fv = "1"}} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %0 = flow.ex.stream.fragment(%arg1 = %cst : vector<3xi32>, %arg2 = %arg0 : tensor<4xf32>) -> tensor<4xf32> { | |
| %1 = flow.dispatch @hloElementwiseOps_ex_dispatch_0::@hloElementwiseOps_rgn_dispatch_0[%arg1 : vector<3xi32>](%arg2) : (tensor<4xf32>) -> tensor<4xf32> | |
| flow.return %1 : tensor<4xf32> | |
| } | |
| return %0 : tensor<4xf32> | |
| } | |
| } | |
| *** IR Dump After iree_compiler::(anonymous namespace)::ConvertFlowToHALPass *** | |
| module @hal_usage { | |
| hal.executable @hloElementwiseOps_ex_dispatch_0 { | |
| hal.executable.entry_point @hloElementwiseOps_rgn_dispatch_0 attributes {ordinal = 0 : i32, workgroup_size = dense<[32, 1, 1]> : vector<3xi32>} | |
| hal.executable.binary attributes {data = dense<[16, 0, 0, 0, 69, 77, 79, 68, 8, 0, 14, 0, 8, 0, 4, 0, 8, 0, 0, 0, 20, 0, 0, 0, 52, 0, 0, 0, 0, 0, 10, 0, 16, 0, 12, 0, 8, 0, 4, 0, 10, 0, 0, 0, 12, 0, 0, 0, 16, 0, 0, 0, 16, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 28, 0, 0, 0, 6, 0, 0, 0, 109, 111, 100, 117, 108, 101, 0, 0, 12, 0, 16, 0, 12, 0, 8, 0, 0, 0, 4, 0, 12, 0, 0, 0, -52, 0, 0, 0, 48, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 126, -1, -1, -1, 4, 0, 0, 0, 2, 0, 0, 0, 68, 0, 0, 0, 4, 0, 0, 0, -52, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -112, -1, -1, -1, 8, 0, 0, 0, 12, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, -20, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -66, -1, -1, -1, 32, 0, 0, 0, 8, 0, 12, 0, 11, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 1, -52, -1, -1, -1, 8, 0, 0, 0, 20, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, 8, 0, 10, 0, 9, 0, 4, 0, 8, 0, 0, 0, 12, 0, 0, 0, 0, 1, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 32, 0, 0, 0, 8, 0, 12, 0, 8, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 7, 0, 0, 0, 112, 0, 0, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 2, 0, 113, 0, 0, 0, 0, 2, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 3, 0, 115, 2, 0, 0, 0, 3, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 4, 0, 119, 3, 0, 0, 0, 4, 0, 0, 3, 1, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 5, 0, 0, 3, 1, 1, 0, 0, 0, 1, 4, 0, 0, 0, 0, 0, 0, 0, 6, 0, 54, 4, 0, 5, 0, 1, 0, 5, 0, 6, 0, 4, 0]> : vector<436xi8>, format = 1230128453 : i32} { | |
| module { | |
| func @hloElementwiseOps_rgn_dispatch_0(%arg0: memref<4xf32>, %arg1: memref<4xf32>) attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : vector<3xi32>, iree.executable.workload = dense<[4, 1, 1]> : vector<3xi32>, iree.ordinal = 0 : i32} { | |
| %0 = "iree_ll_interp.alloc_heap"() : () -> memref<4xf32> | |
| "iree_ll_interp.add_f"(%arg0, %arg0, %0) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> () | |
| %1 = "iree_ll_interp.alloc_heap"() : () -> memref<4xf32> | |
| "iree_ll_interp.sub_f"(%0, %arg0, %1) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> () | |
| %2 = "iree_ll_interp.alloc_heap"() : () -> memref<4xf32> | |
| "iree_ll_interp.mul_f"(%1, %arg0, %2) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> () | |
| %3 = "iree_ll_interp.constant"() {value = dense<0> : tensor<1xi64>} : () -> memref<1xi64> | |
| %4 = "iree_ll_interp.constant"() {value = dense<4> : tensor<1xi64>} : () -> memref<1xi64> | |
| "iree_ll_interp.dynamic_copy"(%2, %3, %arg1, %3, %4) : (memref<4xf32>, memref<1xi64>, memref<4xf32>, memref<1xi64>, memref<1xi64>) -> () | |
| iree.return | |
| } | |
| } | |
| } | |
| hal.executable.binary attributes {data = dense<[8, 0, 0, 0, 83, 80, 86, 69, -92, -1, -1, -1, 104, 4, 0, 0, 112, 0, 0, 0, 4, 0, 0, 0, -18, -1, -1, -1, 4, 0, 0, 0, 1, 0, 0, 0, 12, 0, 0, 0, 0, 0, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 52, 0, 0, 0, 16, 0, 0, 0, 12, 0, 20, 0, 4, 0, 8, 0, 12, 0, 16, 0, 12, 0, 0, 0, 1, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, 12, 0, 16, 0, 0, 0, 4, 0, 8, 0, 12, 0, 12, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, -4, 0, 0, 0, 3, 2, 35, 7, 0, 0, 1, 0, 22, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 17, 0, 2, 0, 1, 0, 0, 0, 10, 0, 11, 0, 83, 80, 86, 95, 75, 72, 82, 95, 115, 116, 111, 114, 97, 103, 101, 95, 98, 117, 102, 102, 101, 114, 95, 115, 116, 111, 114, 97, 103, 101, 95, 99, 108, 97, 115, 115, 0, 0, 0, 0, 14, 0, 3, 0, 0, 0, 0, 0, 1, 0, 0, 0, 15, 0, 13, 0, 5, 0, 0, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 4, 0, 0, 0, 16, 0, 6, 0, 14, 0, 0, 0, 17, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 5, 0, 7, 0, 4, 0, 0, 0, 103, 108, 111, 98, 97, 108, 73, 110, 118, 111, 99, 97, 116, 105, 111, 110, 73, 68, 0, 0, 5, 0, 12, 0, 10, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 48, 0, 0, 5, 0, 12, 0, 11, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 49, 0, 0, 5, 0, 11, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 71, 0, 4, 0, 4, 0, 0, 0, 11, 0, 0, 0, 28, 0, 0, 0, 71, 0, 4, 0, 7, 0, 0, 0, 6, 0, 0, 0, 4, 0, 0, 0, 72, 0, 5, 0, 6, 0, 0, 0, 0, 0, 0, 0, 35, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 33, 0, 0, 0, 1, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 21, 0, 4, 0, 3, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 23, 0, 4, 0, 2, 0, 0, 0, 3, 0, 0, 0, 3, 0, 0, 0, 32, 0, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 2, 0, 0, 0, 59, 0, 4, 0, 1, 0, 0, 0, 4, 0, 0, 0, 1, 0, 0, 0, 22, 0, 3, 0, 8, 0, 0, 0, 32, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 9, 0, 0, 0, 4, 0, 0, 0, 28, 0, 4, 0, 7, 0, 0, 0, 8, 0, 0, 0, 9, 0, 0, 0, 30, 0, 3, 0, 6, 0, 0, 0, 7, 0, 0, 0, 32, 0, 4, 0, 5, 0, 0, 0, 12, 0, 0, 0, 6, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 10, 0, 0, 0, 12, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 11, 0, 0, 0, 12, 0, 0, 0, 19, 0, 2, 0, 13, 0, 0, 0, 33, 0, 3, 0, 12, 0, 0, 0, 13, 0, 0, 0, 20, 0, 2, 0, 16, 0, 0, 0, 41, 0, 3, 0, 16, 0, 0, 0, 17, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 18, 0, 0, 0, 0, 0, 0, 0, 32, 0, 4, 0, 26, 0, 0, 0, 12, 0, 0, 0, 8, 0, 0, 0, 54, 0, 5, 0, 13, 0, 0, 0, 14, 0, 0, 0, 0, 0, 0, 0, 12, 0, 0, 0, -8, 0, 2, 0, 15, 0, 0, 0, 61, 0, 4, 0, 2, 0, 0, 0, 19, 0, 0, 0, 4, 0, 0, 0, 81, 0, 5, 0, 3, 0, 0, 0, 20, 0, 0, 0, 19, 0, 0, 0, 0, 0, 0, 0, -79, 0, 5, 0, 16, 0, 0, 0, 21, 0, 0, 0, 20, 0, 0, 0, 9, 0, 0, 0, -89, 0, 5, 0, 16, 0, 0, 0, 22, 0, 0, 0, 21, 0, 0, 0, 17, 0, 0, 0, -9, 0, 3, 0, 25, 0, 0, 0, 0, 0, 0, 0, -6, 0, 4, 0, 22, 0, 0, 0, 24, 0, 0, 0, 25, 0, 0, 0, -8, 0, 2, 0, 24, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 27, 0, 0, 0, 10, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 61, 0, 4, 0, 8, 0, 0, 0, 28, 0, 0, 0, 27, 0, 0, 0, -127, 0, 5, 0, 8, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, 28, 0, 0, 0, -125, 0, 5, 0, 8, 0, 0, 0, 30, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, -123, 0, 5, 0, 8, 0, 0, 0, 31, 0, 0, 0, 30, 0, 0, 0, 28, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 32, 0, 0, 0, 11, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 62, 0, 3, 0, 32, 0, 0, 0, 31, 0, 0, 0, -7, 0, 2, 0, 25, 0, 0, 0, -8, 0, 2, 0, 25, 0, 0, 0, -3, 0, 1, 0, 56, 0, 1, 0, 1, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0]> : vector<1188xi8>, format = 1397773893 : i32} { | |
| module { | |
| spv.module "Logical" "GLSL450" { | |
| spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> | |
| spv.globalVariable @hloElementwiseOps_rgn_dispatch_0_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| spv.globalVariable @hloElementwiseOps_rgn_dispatch_0_arg_1 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| func @hloElementwiseOps_rgn_dispatch_0() { | |
| %0 = spv.constant true | |
| %1 = spv.constant 4 : i32 | |
| %2 = spv.constant 0 : i32 | |
| %3 = spv._address_of @hloElementwiseOps_rgn_dispatch_0_arg_1 : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| %4 = spv._address_of @hloElementwiseOps_rgn_dispatch_0_arg_0 : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| %5 = spv._address_of @globalInvocationID : !spv.ptr<vector<3xi32>, Input> | |
| %6 = spv.Load "Input" %5 : vector<3xi32> | |
| %7 = spv.CompositeExtract %6[0 : i32] : vector<3xi32> | |
| %8 = spv.SLessThan %7, %1 : i32 | |
| %9 = spv.LogicalAnd %8, %0 : i1 | |
| spv.selection { | |
| spv.BranchConditional %9, ^bb1, ^bb2 | |
| ^bb1: // pred: ^bb0 | |
| %10 = spv.AccessChain %4[%2, %7] : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| %11 = spv.Load "StorageBuffer" %10 : f32 | |
| %12 = spv.FAdd %11, %11 : f32 | |
| %13 = spv.FSub %12, %11 : f32 | |
| %14 = spv.FMul %13, %11 : f32 | |
| %15 = spv.AccessChain %3[%2, %7] : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| spv.Store "StorageBuffer" %15, %14 : f32 | |
| spv.Branch ^bb2 | |
| ^bb2: // 2 preds: ^bb0, ^bb1 | |
| spv._merge | |
| } | |
| spv.Return | |
| } | |
| spv.EntryPoint "GLCompute" @hloElementwiseOps_rgn_dispatch_0, @globalInvocationID | |
| spv.ExecutionMode @hloElementwiseOps_rgn_dispatch_0 "LocalSize", 32, 1, 1 | |
| } attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]} | |
| } | |
| } | |
| } | |
| func @hloElementwiseOps(%arg0: !iree.ref<!hal.buffer>) -> !iree.ref<!hal.buffer> attributes {iree.module.export, iree.reflection = {f = "I6!B3!d4R6!B3!d4", fv = "1"}} { | |
| %cst = constant dense<[4, 1, 1]> : vector<3xi32> | |
| %dev = hal.ex.shared_device : !iree.ref<!hal.device> | |
| %allocator = hal.device.allocator %dev : !iree.ref<!hal.allocator> | |
| %c4_i32 = constant 4 : i32 | |
| %sz = hal.allocator.compute_size %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", shape=[%c4_i32], element_size=4 | |
| %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !iree.ref<!hal.buffer> | |
| hal.ex.defer_release %buffer : !iree.ref<!hal.buffer> | |
| %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !iree.ref<!hal.command_buffer> | |
| hal.command_buffer.begin %cmd | |
| %exe = hal.ex.cache_executable %dev, @hloElementwiseOps_ex_dispatch_0 : !iree.ref<!hal.executable> | |
| %c0 = constant 0 : index | |
| %c4_i32_0 = constant 4 : i32 | |
| %c32_i32 = constant 32 : i32 | |
| %c36_i32 = constant 36 : i32 | |
| %c1_i32 = constant 1 : i32 | |
| %c35_i32 = constant 35 : i32 | |
| %c1_i32_1 = constant 1 : i32 | |
| %c1 = constant 1 : index | |
| %c1_i32_2 = constant 1 : i32 | |
| %c1_i32_3 = constant 1 : i32 | |
| %c2_i32 = constant 2 : i32 | |
| %c1_i32_4 = constant 1 : i32 | |
| %c1_i32_5 = constant 1 : i32 | |
| %c1_i32_6 = constant 1 : i32 | |
| %c2 = constant 2 : index | |
| %c1_i32_7 = constant 1 : i32 | |
| %c1_i32_8 = constant 1 : i32 | |
| %c2_i32_9 = constant 2 : i32 | |
| %c1_i32_10 = constant 1 : i32 | |
| %c1_i32_11 = constant 1 : i32 | |
| %c1_i32_12 = constant 1 : i32 | |
| %c4_i32_13 = constant 4 : i32 | |
| hal.ex.push_binding %cmd, 0, %arg0, shape=[%c4_i32_13], element_size=4 | |
| hal.ex.defer_release %arg0 : !iree.ref<!hal.buffer> | |
| %c4_i32_14 = constant 4 : i32 | |
| hal.ex.push_binding %cmd, 1, %buffer, shape=[%c4_i32_14], element_size=4 | |
| hal.ex.defer_release %buffer : !iree.ref<!hal.buffer> | |
| hal.command_buffer.dispatch %cmd, %exe, entry_point=0, workgroup_xyz=[%c1_i32_1, %c1_i32_6, %c1_i32_12] | |
| %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32> | |
| hal.command_buffer.execution_barrier %cmd, "CommandRetire", "CommandIssue", memory_barriers=[%memory_barrier] | |
| hal.command_buffer.end %cmd | |
| hal.ex.submit_and_wait %dev, %cmd | |
| return %buffer : !iree.ref<!hal.buffer> | |
| } | |
| } | |
| *** IR Dump After Canonicalizer *** | |
| func @hloElementwiseOps(%arg0: !iree.ref<!hal.buffer>) -> !iree.ref<!hal.buffer> attributes {iree.module.export, iree.reflection = {f = "I6!B3!d4R6!B3!d4", fv = "1"}} { | |
| %c1_i32 = constant 1 : i32 | |
| %c4_i32 = constant 4 : i32 | |
| %dev = hal.ex.shared_device : !iree.ref<!hal.device> | |
| %allocator = hal.device.allocator %dev : !iree.ref<!hal.allocator> | |
| %buffer = hal.allocator.allocate.shaped %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", shape=[%c4_i32], element_size=4 : !iree.ref<!hal.buffer> | |
| hal.ex.defer_release %buffer : !iree.ref<!hal.buffer> | |
| %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !iree.ref<!hal.command_buffer> | |
| hal.command_buffer.begin %cmd | |
| %exe = hal.ex.cache_executable %dev, @hloElementwiseOps_ex_dispatch_0 : !iree.ref<!hal.executable> | |
| hal.ex.push_binding %cmd, 0, %arg0, shape=[%c4_i32], element_size=4 | |
| hal.ex.defer_release %arg0 : !iree.ref<!hal.buffer> | |
| hal.ex.push_binding %cmd, 1, %buffer, shape=[%c4_i32], element_size=4 | |
| hal.ex.defer_release %buffer : !iree.ref<!hal.buffer> | |
| hal.command_buffer.dispatch %cmd, %exe, entry_point=0, workgroup_xyz=[%c1_i32, %c1_i32, %c1_i32] | |
| %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32> | |
| hal.command_buffer.execution_barrier %cmd, "CommandRetire", "CommandIssue", memory_barriers=[%memory_barrier] | |
| hal.command_buffer.end %cmd | |
| hal.ex.submit_and_wait %dev, %cmd | |
| return %buffer : !iree.ref<!hal.buffer> | |
| } | |
| *** IR Dump After CSE *** | |
| func @hloElementwiseOps(%arg0: !iree.ref<!hal.buffer>) -> !iree.ref<!hal.buffer> attributes {iree.module.export, iree.reflection = {f = "I6!B3!d4R6!B3!d4", fv = "1"}} { | |
| %c1_i32 = constant 1 : i32 | |
| %c4_i32 = constant 4 : i32 | |
| %dev = hal.ex.shared_device : !iree.ref<!hal.device> | |
| %allocator = hal.device.allocator %dev : !iree.ref<!hal.allocator> | |
| %buffer = hal.allocator.allocate.shaped %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", shape=[%c4_i32], element_size=4 : !iree.ref<!hal.buffer> | |
| hal.ex.defer_release %buffer : !iree.ref<!hal.buffer> | |
| %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !iree.ref<!hal.command_buffer> | |
| hal.command_buffer.begin %cmd | |
| %exe = hal.ex.cache_executable %dev, @hloElementwiseOps_ex_dispatch_0 : !iree.ref<!hal.executable> | |
| hal.ex.push_binding %cmd, 0, %arg0, shape=[%c4_i32], element_size=4 | |
| hal.ex.defer_release %arg0 : !iree.ref<!hal.buffer> | |
| hal.ex.push_binding %cmd, 1, %buffer, shape=[%c4_i32], element_size=4 | |
| hal.ex.defer_release %buffer : !iree.ref<!hal.buffer> | |
| hal.command_buffer.dispatch %cmd, %exe, entry_point=0, workgroup_xyz=[%c1_i32, %c1_i32, %c1_i32] | |
| %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32> | |
| hal.command_buffer.execution_barrier %cmd, "CommandRetire", "CommandIssue", memory_barriers=[%memory_barrier] | |
| hal.command_buffer.end %cmd | |
| hal.ex.submit_and_wait %dev, %cmd | |
| return %buffer : !iree.ref<!hal.buffer> | |
| } | |
| *** IR Dump After Canonicalizer *** | |
| module @hal_usage { | |
| hal.executable @hloElementwiseOps_ex_dispatch_0 { | |
| hal.executable.entry_point @hloElementwiseOps_rgn_dispatch_0 attributes {ordinal = 0 : i32, workgroup_size = dense<[32, 1, 1]> : vector<3xi32>} | |
| hal.executable.binary attributes {data = dense<[16, 0, 0, 0, 69, 77, 79, 68, 8, 0, 14, 0, 8, 0, 4, 0, 8, 0, 0, 0, 20, 0, 0, 0, 52, 0, 0, 0, 0, 0, 10, 0, 16, 0, 12, 0, 8, 0, 4, 0, 10, 0, 0, 0, 12, 0, 0, 0, 16, 0, 0, 0, 16, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 28, 0, 0, 0, 6, 0, 0, 0, 109, 111, 100, 117, 108, 101, 0, 0, 12, 0, 16, 0, 12, 0, 8, 0, 0, 0, 4, 0, 12, 0, 0, 0, -52, 0, 0, 0, 48, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 126, -1, -1, -1, 4, 0, 0, 0, 2, 0, 0, 0, 68, 0, 0, 0, 4, 0, 0, 0, -52, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -112, -1, -1, -1, 8, 0, 0, 0, 12, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, -20, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -66, -1, -1, -1, 32, 0, 0, 0, 8, 0, 12, 0, 11, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 1, -52, -1, -1, -1, 8, 0, 0, 0, 20, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, 8, 0, 10, 0, 9, 0, 4, 0, 8, 0, 0, 0, 12, 0, 0, 0, 0, 1, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 32, 0, 0, 0, 8, 0, 12, 0, 8, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 7, 0, 0, 0, 112, 0, 0, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 2, 0, 113, 0, 0, 0, 0, 2, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 3, 0, 115, 2, 0, 0, 0, 3, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 4, 0, 119, 3, 0, 0, 0, 4, 0, 0, 3, 1, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 5, 0, 0, 3, 1, 1, 0, 0, 0, 1, 4, 0, 0, 0, 0, 0, 0, 0, 6, 0, 54, 4, 0, 5, 0, 1, 0, 5, 0, 6, 0, 4, 0]> : vector<436xi8>, format = 1230128453 : i32} { | |
| module { | |
| func @hloElementwiseOps_rgn_dispatch_0(%arg0: memref<4xf32>, %arg1: memref<4xf32>) attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : vector<3xi32>, iree.executable.workload = dense<[4, 1, 1]> : vector<3xi32>, iree.ordinal = 0 : i32} { | |
| %0 = "iree_ll_interp.alloc_heap"() : () -> memref<4xf32> | |
| "iree_ll_interp.add_f"(%arg0, %arg0, %0) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> () | |
| %1 = "iree_ll_interp.alloc_heap"() : () -> memref<4xf32> | |
| "iree_ll_interp.sub_f"(%0, %arg0, %1) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> () | |
| %2 = "iree_ll_interp.alloc_heap"() : () -> memref<4xf32> | |
| "iree_ll_interp.mul_f"(%1, %arg0, %2) : (memref<4xf32>, memref<4xf32>, memref<4xf32>) -> () | |
| %3 = "iree_ll_interp.constant"() {value = dense<0> : tensor<1xi64>} : () -> memref<1xi64> | |
| %4 = "iree_ll_interp.constant"() {value = dense<4> : tensor<1xi64>} : () -> memref<1xi64> | |
| "iree_ll_interp.dynamic_copy"(%2, %3, %arg1, %3, %4) : (memref<4xf32>, memref<1xi64>, memref<4xf32>, memref<1xi64>, memref<1xi64>) -> () | |
| iree.return | |
| } | |
| } | |
| } | |
| hal.executable.binary attributes {data = dense<[8, 0, 0, 0, 83, 80, 86, 69, -92, -1, -1, -1, 104, 4, 0, 0, 112, 0, 0, 0, 4, 0, 0, 0, -18, -1, -1, -1, 4, 0, 0, 0, 1, 0, 0, 0, 12, 0, 0, 0, 0, 0, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 52, 0, 0, 0, 16, 0, 0, 0, 12, 0, 20, 0, 4, 0, 8, 0, 12, 0, 16, 0, 12, 0, 0, 0, 1, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, 12, 0, 16, 0, 0, 0, 4, 0, 8, 0, 12, 0, 12, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, -4, 0, 0, 0, 3, 2, 35, 7, 0, 0, 1, 0, 22, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 17, 0, 2, 0, 1, 0, 0, 0, 10, 0, 11, 0, 83, 80, 86, 95, 75, 72, 82, 95, 115, 116, 111, 114, 97, 103, 101, 95, 98, 117, 102, 102, 101, 114, 95, 115, 116, 111, 114, 97, 103, 101, 95, 99, 108, 97, 115, 115, 0, 0, 0, 0, 14, 0, 3, 0, 0, 0, 0, 0, 1, 0, 0, 0, 15, 0, 13, 0, 5, 0, 0, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 4, 0, 0, 0, 16, 0, 6, 0, 14, 0, 0, 0, 17, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 5, 0, 7, 0, 4, 0, 0, 0, 103, 108, 111, 98, 97, 108, 73, 110, 118, 111, 99, 97, 116, 105, 111, 110, 73, 68, 0, 0, 5, 0, 12, 0, 10, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 48, 0, 0, 5, 0, 12, 0, 11, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 49, 0, 0, 5, 0, 11, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 71, 0, 4, 0, 4, 0, 0, 0, 11, 0, 0, 0, 28, 0, 0, 0, 71, 0, 4, 0, 7, 0, 0, 0, 6, 0, 0, 0, 4, 0, 0, 0, 72, 0, 5, 0, 6, 0, 0, 0, 0, 0, 0, 0, 35, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 33, 0, 0, 0, 1, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 21, 0, 4, 0, 3, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 23, 0, 4, 0, 2, 0, 0, 0, 3, 0, 0, 0, 3, 0, 0, 0, 32, 0, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 2, 0, 0, 0, 59, 0, 4, 0, 1, 0, 0, 0, 4, 0, 0, 0, 1, 0, 0, 0, 22, 0, 3, 0, 8, 0, 0, 0, 32, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 9, 0, 0, 0, 4, 0, 0, 0, 28, 0, 4, 0, 7, 0, 0, 0, 8, 0, 0, 0, 9, 0, 0, 0, 30, 0, 3, 0, 6, 0, 0, 0, 7, 0, 0, 0, 32, 0, 4, 0, 5, 0, 0, 0, 12, 0, 0, 0, 6, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 10, 0, 0, 0, 12, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 11, 0, 0, 0, 12, 0, 0, 0, 19, 0, 2, 0, 13, 0, 0, 0, 33, 0, 3, 0, 12, 0, 0, 0, 13, 0, 0, 0, 20, 0, 2, 0, 16, 0, 0, 0, 41, 0, 3, 0, 16, 0, 0, 0, 17, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 18, 0, 0, 0, 0, 0, 0, 0, 32, 0, 4, 0, 26, 0, 0, 0, 12, 0, 0, 0, 8, 0, 0, 0, 54, 0, 5, 0, 13, 0, 0, 0, 14, 0, 0, 0, 0, 0, 0, 0, 12, 0, 0, 0, -8, 0, 2, 0, 15, 0, 0, 0, 61, 0, 4, 0, 2, 0, 0, 0, 19, 0, 0, 0, 4, 0, 0, 0, 81, 0, 5, 0, 3, 0, 0, 0, 20, 0, 0, 0, 19, 0, 0, 0, 0, 0, 0, 0, -79, 0, 5, 0, 16, 0, 0, 0, 21, 0, 0, 0, 20, 0, 0, 0, 9, 0, 0, 0, -89, 0, 5, 0, 16, 0, 0, 0, 22, 0, 0, 0, 21, 0, 0, 0, 17, 0, 0, 0, -9, 0, 3, 0, 25, 0, 0, 0, 0, 0, 0, 0, -6, 0, 4, 0, 22, 0, 0, 0, 24, 0, 0, 0, 25, 0, 0, 0, -8, 0, 2, 0, 24, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 27, 0, 0, 0, 10, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 61, 0, 4, 0, 8, 0, 0, 0, 28, 0, 0, 0, 27, 0, 0, 0, -127, 0, 5, 0, 8, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, 28, 0, 0, 0, -125, 0, 5, 0, 8, 0, 0, 0, 30, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, -123, 0, 5, 0, 8, 0, 0, 0, 31, 0, 0, 0, 30, 0, 0, 0, 28, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 32, 0, 0, 0, 11, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 62, 0, 3, 0, 32, 0, 0, 0, 31, 0, 0, 0, -7, 0, 2, 0, 25, 0, 0, 0, -8, 0, 2, 0, 25, 0, 0, 0, -3, 0, 1, 0, 56, 0, 1, 0, 1, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0]> : vector<1188xi8>, format = 1397773893 : i32} { | |
| module { | |
| spv.module "Logical" "GLSL450" { | |
| spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> | |
| spv.globalVariable @hloElementwiseOps_rgn_dispatch_0_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| spv.globalVariable @hloElementwiseOps_rgn_dispatch_0_arg_1 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| func @hloElementwiseOps_rgn_dispatch_0() { | |
| %0 = spv.constant true | |
| %1 = spv.constant 4 : i32 | |
| %2 = spv.constant 0 : i32 | |
| %3 = spv._address_of @hloElementwiseOps_rgn_dispatch_0_arg_1 : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| %4 = spv._address_of @hloElementwiseOps_rgn_dispatch_0_arg_0 : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| %5 = spv._address_of @globalInvocationID : !spv.ptr<vector<3xi32>, Input> | |
| %6 = spv.Load "Input" %5 : vector<3xi32> | |
| %7 = spv.CompositeExtract %6[0 : i32] : vector<3xi32> | |
| %8 = spv.SLessThan %7, %1 : i32 | |
| %9 = spv.LogicalAnd %8, %0 : i1 | |
| spv.selection { | |
| spv.BranchConditional %9, ^bb1, ^bb2 | |
| ^bb1: // pred: ^bb0 | |
| %10 = spv.AccessChain %4[%2, %7] : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| %11 = spv.Load "StorageBuffer" %10 : f32 | |
| %12 = spv.FAdd %11, %11 : f32 | |
| %13 = spv.FSub %12, %11 : f32 | |
| %14 = spv.FMul %13, %11 : f32 | |
| %15 = spv.AccessChain %3[%2, %7] : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> | |
| spv.Store "StorageBuffer" %15, %14 : f32 | |
| spv.Branch ^bb2 | |
| ^bb2: // 2 preds: ^bb0, ^bb1 | |
| spv._merge | |
| } | |
| spv.Return | |
| } | |
| spv.EntryPoint "GLCompute" @hloElementwiseOps_rgn_dispatch_0, @globalInvocationID | |
| spv.ExecutionMode @hloElementwiseOps_rgn_dispatch_0 "LocalSize", 32, 1, 1 | |
| } attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]} | |
| } | |
| } | |
| } | |
| func @hloElementwiseOps(%arg0: !iree.ref<!hal.buffer>) -> !iree.ref<!hal.buffer> attributes {iree.module.export, iree.reflection = {f = "I6!B3!d4R6!B3!d4", fv = "1"}} { | |
| %c1_i32 = constant 1 : i32 | |
| %c4_i32 = constant 4 : i32 | |
| %dev = hal.ex.shared_device : !iree.ref<!hal.device> | |
| %allocator = hal.device.allocator %dev : !iree.ref<!hal.allocator> | |
| %buffer = hal.allocator.allocate.shaped %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", shape=[%c4_i32], element_size=4 : !iree.ref<!hal.buffer> | |
| hal.ex.defer_release %buffer : !iree.ref<!hal.buffer> | |
| %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !iree.ref<!hal.command_buffer> | |
| hal.command_buffer.begin %cmd | |
| %exe = hal.ex.cache_executable %dev, @hloElementwiseOps_ex_dispatch_0 : !iree.ref<!hal.executable> | |
| hal.ex.push_binding %cmd, 0, %arg0, shape=[%c4_i32], element_size=4 | |
| hal.ex.defer_release %arg0 : !iree.ref<!hal.buffer> | |
| hal.ex.push_binding %cmd, 1, %buffer, shape=[%c4_i32], element_size=4 | |
| hal.ex.defer_release %buffer : !iree.ref<!hal.buffer> | |
| hal.command_buffer.dispatch %cmd, %exe, entry_point=0, workgroup_xyz=[%c1_i32, %c1_i32, %c1_i32] | |
| %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32> | |
| hal.command_buffer.execution_barrier %cmd, "CommandRetire", "CommandIssue", memory_barriers=[%memory_barrier] | |
| hal.command_buffer.end %cmd | |
| hal.ex.submit_and_wait %dev, %cmd | |
| return %buffer : !iree.ref<!hal.buffer> | |
| } | |
| } | |
| *** IR Dump After iree_compiler::IREE::VM::ConversionPass *** | |
| module @hal_usage { | |
| vm.module @hal_usage { | |
| vm.rodata @hloElementwiseOps_ex_dispatch_0_data_1230128453 dense<[16, 0, 0, 0, 69, 77, 79, 68, 8, 0, 14, 0, 8, 0, 4, 0, 8, 0, 0, 0, 20, 0, 0, 0, 52, 0, 0, 0, 0, 0, 10, 0, 16, 0, 12, 0, 8, 0, 4, 0, 10, 0, 0, 0, 12, 0, 0, 0, 16, 0, 0, 0, 16, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 28, 0, 0, 0, 6, 0, 0, 0, 109, 111, 100, 117, 108, 101, 0, 0, 12, 0, 16, 0, 12, 0, 8, 0, 0, 0, 4, 0, 12, 0, 0, 0, -52, 0, 0, 0, 48, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 126, -1, -1, -1, 4, 0, 0, 0, 2, 0, 0, 0, 68, 0, 0, 0, 4, 0, 0, 0, -52, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -112, -1, -1, -1, 8, 0, 0, 0, 12, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, -20, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -66, -1, -1, -1, 32, 0, 0, 0, 8, 0, 12, 0, 11, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 1, -52, -1, -1, -1, 8, 0, 0, 0, 20, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, 8, 0, 10, 0, 9, 0, 4, 0, 8, 0, 0, 0, 12, 0, 0, 0, 0, 1, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 32, 0, 0, 0, 8, 0, 12, 0, 8, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 7, 0, 0, 0, 112, 0, 0, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 2, 0, 113, 0, 0, 0, 0, 2, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 3, 0, 115, 2, 0, 0, 0, 3, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 4, 0, 119, 3, 0, 0, 0, 4, 0, 0, 3, 1, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 5, 0, 0, 3, 1, 1, 0, 0, 0, 1, 4, 0, 0, 0, 0, 0, 0, 0, 6, 0, 54, 4, 0, 5, 0, 1, 0, 5, 0, 6, 0, 4, 0]> : vector<436xi8> | |
| vm.rodata @hloElementwiseOps_ex_dispatch_0_data_1397773893 dense<[8, 0, 0, 0, 83, 80, 86, 69, -92, -1, -1, -1, 104, 4, 0, 0, 112, 0, 0, 0, 4, 0, 0, 0, -18, -1, -1, -1, 4, 0, 0, 0, 1, 0, 0, 0, 12, 0, 0, 0, 0, 0, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 52, 0, 0, 0, 16, 0, 0, 0, 12, 0, 20, 0, 4, 0, 8, 0, 12, 0, 16, 0, 12, 0, 0, 0, 1, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, 12, 0, 16, 0, 0, 0, 4, 0, 8, 0, 12, 0, 12, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, -4, 0, 0, 0, 3, 2, 35, 7, 0, 0, 1, 0, 22, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 17, 0, 2, 0, 1, 0, 0, 0, 10, 0, 11, 0, 83, 80, 86, 95, 75, 72, 82, 95, 115, 116, 111, 114, 97, 103, 101, 95, 98, 117, 102, 102, 101, 114, 95, 115, 116, 111, 114, 97, 103, 101, 95, 99, 108, 97, 115, 115, 0, 0, 0, 0, 14, 0, 3, 0, 0, 0, 0, 0, 1, 0, 0, 0, 15, 0, 13, 0, 5, 0, 0, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 4, 0, 0, 0, 16, 0, 6, 0, 14, 0, 0, 0, 17, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 5, 0, 7, 0, 4, 0, 0, 0, 103, 108, 111, 98, 97, 108, 73, 110, 118, 111, 99, 97, 116, 105, 111, 110, 73, 68, 0, 0, 5, 0, 12, 0, 10, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 48, 0, 0, 5, 0, 12, 0, 11, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 49, 0, 0, 5, 0, 11, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 71, 0, 4, 0, 4, 0, 0, 0, 11, 0, 0, 0, 28, 0, 0, 0, 71, 0, 4, 0, 7, 0, 0, 0, 6, 0, 0, 0, 4, 0, 0, 0, 72, 0, 5, 0, 6, 0, 0, 0, 0, 0, 0, 0, 35, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 33, 0, 0, 0, 1, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 21, 0, 4, 0, 3, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 23, 0, 4, 0, 2, 0, 0, 0, 3, 0, 0, 0, 3, 0, 0, 0, 32, 0, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 2, 0, 0, 0, 59, 0, 4, 0, 1, 0, 0, 0, 4, 0, 0, 0, 1, 0, 0, 0, 22, 0, 3, 0, 8, 0, 0, 0, 32, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 9, 0, 0, 0, 4, 0, 0, 0, 28, 0, 4, 0, 7, 0, 0, 0, 8, 0, 0, 0, 9, 0, 0, 0, 30, 0, 3, 0, 6, 0, 0, 0, 7, 0, 0, 0, 32, 0, 4, 0, 5, 0, 0, 0, 12, 0, 0, 0, 6, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 10, 0, 0, 0, 12, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 11, 0, 0, 0, 12, 0, 0, 0, 19, 0, 2, 0, 13, 0, 0, 0, 33, 0, 3, 0, 12, 0, 0, 0, 13, 0, 0, 0, 20, 0, 2, 0, 16, 0, 0, 0, 41, 0, 3, 0, 16, 0, 0, 0, 17, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 18, 0, 0, 0, 0, 0, 0, 0, 32, 0, 4, 0, 26, 0, 0, 0, 12, 0, 0, 0, 8, 0, 0, 0, 54, 0, 5, 0, 13, 0, 0, 0, 14, 0, 0, 0, 0, 0, 0, 0, 12, 0, 0, 0, -8, 0, 2, 0, 15, 0, 0, 0, 61, 0, 4, 0, 2, 0, 0, 0, 19, 0, 0, 0, 4, 0, 0, 0, 81, 0, 5, 0, 3, 0, 0, 0, 20, 0, 0, 0, 19, 0, 0, 0, 0, 0, 0, 0, -79, 0, 5, 0, 16, 0, 0, 0, 21, 0, 0, 0, 20, 0, 0, 0, 9, 0, 0, 0, -89, 0, 5, 0, 16, 0, 0, 0, 22, 0, 0, 0, 21, 0, 0, 0, 17, 0, 0, 0, -9, 0, 3, 0, 25, 0, 0, 0, 0, 0, 0, 0, -6, 0, 4, 0, 22, 0, 0, 0, 24, 0, 0, 0, 25, 0, 0, 0, -8, 0, 2, 0, 24, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 27, 0, 0, 0, 10, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 61, 0, 4, 0, 8, 0, 0, 0, 28, 0, 0, 0, 27, 0, 0, 0, -127, 0, 5, 0, 8, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, 28, 0, 0, 0, -125, 0, 5, 0, 8, 0, 0, 0, 30, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, -123, 0, 5, 0, 8, 0, 0, 0, 31, 0, 0, 0, 30, 0, 0, 0, 28, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 32, 0, 0, 0, 11, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 62, 0, 3, 0, 32, 0, 0, 0, 31, 0, 0, 0, -7, 0, 2, 0, 25, 0, 0, 0, -8, 0, 2, 0, 25, 0, 0, 0, -3, 0, 1, 0, 56, 0, 1, 0, 1, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0]> : vector<1188xi8> | |
| vm.global.ref @hloElementwiseOps_ex_dispatch_0_cached mutable : !iree.ref<!hal.executable> | |
| vm.func @hloElementwiseOps_ex_dispatch_0(%arg0: !iree.ref<!hal.device>) -> !iree.ref<!hal.executable> { | |
| %hloElementwiseOps_ex_dispatch_0_cached = vm.global.load.ref @hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable> | |
| %rnz = vm.cmp.nz.ref %hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable> | |
| vm.cond_br %rnz, ^bb1(%hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable>), ^bb2 | |
| ^bb1(%0: !iree.ref<!hal.executable>): // pred: ^bb0 | |
| vm.return %0 : !iree.ref<!hal.executable> | |
| ^bb2: // pred: ^bb0 | |
| %c1230128453 = vm.const.i32 1230128453 : i32 | |
| %c1397773893 = vm.const.i32 1397773893 : i32 | |
| %1 = vm.call.variadic @hal.ex.match_supported_executable_format(%arg0, [%c1230128453, %c1397773893]) : (!iree.ref<!hal.device>, i32...) -> i32 | |
| vm.br ^bb3(%1 : i32) | |
| ^bb3(%2: i32): // pred: ^bb2 | |
| %c1230128453_0 = vm.const.i32 1230128453 : i32 | |
| %eq = vm.cmp.eq.i32 %2, %c1230128453_0 : i32 | |
| vm.cond_br %eq, ^bb4(%2 : i32), ^bb5(%2 : i32) | |
| ^bb4(%3: i32): // pred: ^bb3 | |
| %hloElementwiseOps_ex_dispatch_0_data_1230128453 = vm.const.ref.rodata @hloElementwiseOps_ex_dispatch_0_data_1230128453 : !iree.byte_buffer_ref | |
| %ref = vm.call @hal.ex.cache_executable(%arg0, %3, %hloElementwiseOps_ex_dispatch_0_data_1230128453) : (!iree.ref<!hal.device>, i32, !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> | |
| vm.br ^bb7(%ref : !iree.ref<!hal.executable>) | |
| ^bb5(%4: i32): // pred: ^bb3 | |
| %c1397773893_1 = vm.const.i32 1397773893 : i32 | |
| %eq_2 = vm.cmp.eq.i32 %4, %c1397773893_1 : i32 | |
| vm.cond_br %eq_2, ^bb6(%4 : i32), ^bb8 | |
| ^bb6(%5: i32): // pred: ^bb5 | |
| %hloElementwiseOps_ex_dispatch_0_data_1397773893 = vm.const.ref.rodata @hloElementwiseOps_ex_dispatch_0_data_1397773893 : !iree.byte_buffer_ref | |
| %ref_3 = vm.call @hal.ex.cache_executable(%arg0, %5, %hloElementwiseOps_ex_dispatch_0_data_1397773893) : (!iree.ref<!hal.device>, i32, !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> | |
| vm.br ^bb7(%ref_3 : !iree.ref<!hal.executable>) | |
| ^bb7(%6: !iree.ref<!hal.executable>): // 2 preds: ^bb4, ^bb6 | |
| vm.global.store.ref @hloElementwiseOps_ex_dispatch_0_cached, %6 : !iree.ref<!hal.executable> | |
| vm.return %6 : !iree.ref<!hal.executable> | |
| ^bb8: // pred: ^bb5 | |
| %null = vm.const.ref.zero : !iree.ref<!hal.executable> | |
| vm.return %null : !iree.ref<!hal.executable> | |
| } | |
| vm.func @hloElementwiseOps(%arg0: !iree.ref<!hal.buffer>) -> !iree.ref<!hal.buffer> attributes {iree.reflection = {f = "I6!B3!d4R6!B3!d4", fv = "1"}} { | |
| %c1 = vm.const.i32 1 : i32 | |
| %c4 = vm.const.i32 4 : i32 | |
| %ref = vm.call @hal.ex.shared_device() : () -> !iree.ref<!hal.device> | |
| %ref_0 = vm.call @hal.device.allocator(%ref) : (!iree.ref<!hal.device>) -> !iree.ref<!hal.allocator> | |
| %c50 = vm.const.i32 50 : i32 | |
| %c15 = vm.const.i32 15 : i32 | |
| %c4_1 = vm.const.i32 4 : i32 | |
| %ref_2 = vm.call.variadic @hal.allocator.allocate.shaped(%ref_0, %c50, %c15, [%c4], %c4_1) : (!iree.ref<!hal.allocator>, i32, i32, i32..., i32) -> !iree.ref<!hal.buffer> | |
| vm.call @hal.ex.defer_release(%ref_2) : (!iree.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) : (!iree.ref<!hal.device>, i32, i32) -> !iree.ref<!hal.command_buffer> | |
| vm.call @hal.command_buffer.begin(%ref_4) : (!iree.ref<!hal.command_buffer>) -> () | |
| %ref_5 = vm.call @hloElementwiseOps_ex_dispatch_0(%ref) : (!iree.ref<!hal.device>) -> !iree.ref<!hal.executable> | |
| %zero = vm.const.i32.zero : i32 | |
| %c4_6 = vm.const.i32 4 : i32 | |
| vm.call.variadic @hal.ex.push_binding(%ref_4, %zero, %arg0, [%c4], %c4_6) : (!iree.ref<!hal.command_buffer>, i32, !iree.ref<!hal.buffer>, i32..., i32) | |
| vm.call @hal.ex.defer_release(%arg0) : (!iree.ref<!hal.buffer>) -> () | |
| %c1_7 = vm.const.i32 1 : i32 | |
| %c4_8 = vm.const.i32 4 : i32 | |
| vm.call.variadic @hal.ex.push_binding(%ref_4, %c1_7, %ref_2, [%c4], %c4_8) : (!iree.ref<!hal.command_buffer>, i32, !iree.ref<!hal.buffer>, i32..., i32) | |
| vm.call @hal.ex.defer_release(%ref_2) : (!iree.ref<!hal.buffer>) -> () | |
| %zero_9 = vm.const.i32.zero : i32 | |
| vm.call @hal.command_buffer.dispatch(%ref_4, %ref_5, %zero_9, %c1, %c1, %c1) : (!iree.ref<!hal.command_buffer>, !iree.ref<!hal.executable>, i32, i32, i32, i32) -> () | |
| %c16 = vm.const.i32 16 : i32 | |
| %c1_10 = vm.const.i32 1 : i32 | |
| %c8 = vm.const.i32 8 : i32 | |
| %c4_11 = vm.const.i32 4 : i32 | |
| vm.call.variadic @hal.command_buffer.execution_barrier(%ref_4, %c16, %c1_10, [%c8], []) : (!iree.ref<!hal.command_buffer>, i32, i32, i32..., i32...) | |
| vm.call @hal.command_buffer.end(%ref_4) : (!iree.ref<!hal.command_buffer>) -> () | |
| vm.call @hal.ex.submit_and_wait(%ref, %ref_4) : (!iree.ref<!hal.device>, !iree.ref<!hal.command_buffer>) -> () | |
| vm.return %ref_2 : !iree.ref<!hal.buffer> | |
| } | |
| vm.export @hloElementwiseOps | |
| vm.import @hal.ex.shared_device() -> !iree.ref<!hal.device> attributes {nosideeffects} | |
| vm.import @hal.ex.match_supported_executable_format(%device : !iree.ref<!hal.device>, %available_formats : i32...) -> i32 attributes {nosideeffects} | |
| vm.import @hal.ex.cache_executable(%device : !iree.ref<!hal.device>, %executable_format : i32, %executable_data : !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> attributes {nosideeffects} | |
| vm.import @hal.ex.push_binding(%command_buffer : !iree.ref<!hal.command_buffer>, %ordinal : i32, %buffer : !iree.ref<!hal.buffer>, %shape : i32..., %element_size : i32) | |
| vm.import @hal.ex.executable_descriptor_set_layout(%executable : !iree.ref<!hal.executable>, %set : i32) -> !iree.ref<!hal.descriptor_set_layout> | |
| vm.import @hal.ex.defer_release(%operand : !iree.opaque_ref) | |
| vm.import @hal.ex.submit_and_wait(%device : !iree.ref<!hal.device>, %command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.allocator.compute_size(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32) -> i32 attributes {nosideeffects} | |
| vm.import @hal.allocator.allocate(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.allocator.allocate.const(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32, %value : !iree.byte_buffer_ref) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.allocator.allocate.shaped(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.buffer.subspan(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.buffer.fill(%target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) | |
| vm.import @hal.buffer.read_data(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.mutable_byte_buffer_ref, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer.write_data(%target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %source_buffer : !iree.byte_buffer_ref, %source_offset : i32, %length : i32) | |
| vm.import @hal.buffer.copy_data(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer.load(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 | |
| vm.import @hal.buffer.store(%value : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer_view.compute_offset(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %element_size : i32) -> i32 | |
| vm.import @hal.buffer_view.compute_length(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %element_size : i32) -> i32 | |
| vm.import @hal.buffer_view.compute_range(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %lengths : i32..., %element_size : i32) -> (i32, i32) | |
| vm.import @hal.buffer_view.slice(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %lengths : i32..., %element_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.command_buffer.create(%device : !iree.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !iree.ref<!hal.command_buffer> | |
| vm.import @hal.command_buffer.begin(%command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.command_buffer.end(%command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.command_buffer.execution_barrier(%command_buffer : !iree.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) | |
| vm.import @hal.command_buffer.fill_buffer(%command_buffer : !iree.ref<!hal.command_buffer>, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) | |
| vm.import @hal.command_buffer.copy_buffer(%command_buffer : !iree.ref<!hal.command_buffer>, %source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.command_buffer.bind_descriptor_set(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %set : i32, %descriptor_set : !iree.ref<!hal.descriptor_set>, %dynamic_offsets : i32...) | |
| vm.import @hal.command_buffer.dispatch(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) | |
| vm.import @hal.command_buffer.dispatch.indirect(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %entry_point : i32, %workgroups_buffer : !iree.ref<!hal.buffer>, %workgroups_offset : i32) | |
| vm.import @hal.descriptor_set.allocate(%device : !iree.ref<!hal.device>, %set_layout : !iree.ref<!hal.descriptor_set_layout>) -> !iree.ref<!hal.descriptor_set> | |
| vm.import @hal.descriptor_set.update(%device : !iree.ref<!hal.device>, %set : !iree.ref<!hal.descriptor_set>, %binding : i32, %buffer : !iree.ref<!hal.buffer>, %offset : i32, %length : i32, %access : i32) | |
| vm.import @hal.device.allocator(%device : !iree.ref<!hal.device>) -> !iree.ref<!hal.allocator> attributes {nosideeffects} | |
| } | |
| } | |
| *** IR Dump After iree_compiler::IREE::VM::GlobalInitializationPass *** | |
| vm.module @hal_usage { | |
| vm.rodata @hloElementwiseOps_ex_dispatch_0_data_1230128453 dense<[16, 0, 0, 0, 69, 77, 79, 68, 8, 0, 14, 0, 8, 0, 4, 0, 8, 0, 0, 0, 20, 0, 0, 0, 52, 0, 0, 0, 0, 0, 10, 0, 16, 0, 12, 0, 8, 0, 4, 0, 10, 0, 0, 0, 12, 0, 0, 0, 16, 0, 0, 0, 16, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 28, 0, 0, 0, 6, 0, 0, 0, 109, 111, 100, 117, 108, 101, 0, 0, 12, 0, 16, 0, 12, 0, 8, 0, 0, 0, 4, 0, 12, 0, 0, 0, -52, 0, 0, 0, 48, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 126, -1, -1, -1, 4, 0, 0, 0, 2, 0, 0, 0, 68, 0, 0, 0, 4, 0, 0, 0, -52, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -112, -1, -1, -1, 8, 0, 0, 0, 12, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, -20, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -66, -1, -1, -1, 32, 0, 0, 0, 8, 0, 12, 0, 11, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 1, -52, -1, -1, -1, 8, 0, 0, 0, 20, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, 8, 0, 10, 0, 9, 0, 4, 0, 8, 0, 0, 0, 12, 0, 0, 0, 0, 1, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 32, 0, 0, 0, 8, 0, 12, 0, 8, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 7, 0, 0, 0, 112, 0, 0, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 2, 0, 113, 0, 0, 0, 0, 2, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 3, 0, 115, 2, 0, 0, 0, 3, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 4, 0, 119, 3, 0, 0, 0, 4, 0, 0, 3, 1, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 5, 0, 0, 3, 1, 1, 0, 0, 0, 1, 4, 0, 0, 0, 0, 0, 0, 0, 6, 0, 54, 4, 0, 5, 0, 1, 0, 5, 0, 6, 0, 4, 0]> : vector<436xi8> | |
| vm.rodata @hloElementwiseOps_ex_dispatch_0_data_1397773893 dense<[8, 0, 0, 0, 83, 80, 86, 69, -92, -1, -1, -1, 104, 4, 0, 0, 112, 0, 0, 0, 4, 0, 0, 0, -18, -1, -1, -1, 4, 0, 0, 0, 1, 0, 0, 0, 12, 0, 0, 0, 0, 0, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 52, 0, 0, 0, 16, 0, 0, 0, 12, 0, 20, 0, 4, 0, 8, 0, 12, 0, 16, 0, 12, 0, 0, 0, 1, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, 12, 0, 16, 0, 0, 0, 4, 0, 8, 0, 12, 0, 12, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, -4, 0, 0, 0, 3, 2, 35, 7, 0, 0, 1, 0, 22, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 17, 0, 2, 0, 1, 0, 0, 0, 10, 0, 11, 0, 83, 80, 86, 95, 75, 72, 82, 95, 115, 116, 111, 114, 97, 103, 101, 95, 98, 117, 102, 102, 101, 114, 95, 115, 116, 111, 114, 97, 103, 101, 95, 99, 108, 97, 115, 115, 0, 0, 0, 0, 14, 0, 3, 0, 0, 0, 0, 0, 1, 0, 0, 0, 15, 0, 13, 0, 5, 0, 0, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 4, 0, 0, 0, 16, 0, 6, 0, 14, 0, 0, 0, 17, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 5, 0, 7, 0, 4, 0, 0, 0, 103, 108, 111, 98, 97, 108, 73, 110, 118, 111, 99, 97, 116, 105, 111, 110, 73, 68, 0, 0, 5, 0, 12, 0, 10, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 48, 0, 0, 5, 0, 12, 0, 11, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 49, 0, 0, 5, 0, 11, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 71, 0, 4, 0, 4, 0, 0, 0, 11, 0, 0, 0, 28, 0, 0, 0, 71, 0, 4, 0, 7, 0, 0, 0, 6, 0, 0, 0, 4, 0, 0, 0, 72, 0, 5, 0, 6, 0, 0, 0, 0, 0, 0, 0, 35, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 33, 0, 0, 0, 1, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 21, 0, 4, 0, 3, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 23, 0, 4, 0, 2, 0, 0, 0, 3, 0, 0, 0, 3, 0, 0, 0, 32, 0, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 2, 0, 0, 0, 59, 0, 4, 0, 1, 0, 0, 0, 4, 0, 0, 0, 1, 0, 0, 0, 22, 0, 3, 0, 8, 0, 0, 0, 32, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 9, 0, 0, 0, 4, 0, 0, 0, 28, 0, 4, 0, 7, 0, 0, 0, 8, 0, 0, 0, 9, 0, 0, 0, 30, 0, 3, 0, 6, 0, 0, 0, 7, 0, 0, 0, 32, 0, 4, 0, 5, 0, 0, 0, 12, 0, 0, 0, 6, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 10, 0, 0, 0, 12, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 11, 0, 0, 0, 12, 0, 0, 0, 19, 0, 2, 0, 13, 0, 0, 0, 33, 0, 3, 0, 12, 0, 0, 0, 13, 0, 0, 0, 20, 0, 2, 0, 16, 0, 0, 0, 41, 0, 3, 0, 16, 0, 0, 0, 17, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 18, 0, 0, 0, 0, 0, 0, 0, 32, 0, 4, 0, 26, 0, 0, 0, 12, 0, 0, 0, 8, 0, 0, 0, 54, 0, 5, 0, 13, 0, 0, 0, 14, 0, 0, 0, 0, 0, 0, 0, 12, 0, 0, 0, -8, 0, 2, 0, 15, 0, 0, 0, 61, 0, 4, 0, 2, 0, 0, 0, 19, 0, 0, 0, 4, 0, 0, 0, 81, 0, 5, 0, 3, 0, 0, 0, 20, 0, 0, 0, 19, 0, 0, 0, 0, 0, 0, 0, -79, 0, 5, 0, 16, 0, 0, 0, 21, 0, 0, 0, 20, 0, 0, 0, 9, 0, 0, 0, -89, 0, 5, 0, 16, 0, 0, 0, 22, 0, 0, 0, 21, 0, 0, 0, 17, 0, 0, 0, -9, 0, 3, 0, 25, 0, 0, 0, 0, 0, 0, 0, -6, 0, 4, 0, 22, 0, 0, 0, 24, 0, 0, 0, 25, 0, 0, 0, -8, 0, 2, 0, 24, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 27, 0, 0, 0, 10, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 61, 0, 4, 0, 8, 0, 0, 0, 28, 0, 0, 0, 27, 0, 0, 0, -127, 0, 5, 0, 8, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, 28, 0, 0, 0, -125, 0, 5, 0, 8, 0, 0, 0, 30, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, -123, 0, 5, 0, 8, 0, 0, 0, 31, 0, 0, 0, 30, 0, 0, 0, 28, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 32, 0, 0, 0, 11, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 62, 0, 3, 0, 32, 0, 0, 0, 31, 0, 0, 0, -7, 0, 2, 0, 25, 0, 0, 0, -8, 0, 2, 0, 25, 0, 0, 0, -3, 0, 1, 0, 56, 0, 1, 0, 1, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0]> : vector<1188xi8> | |
| vm.global.ref @hloElementwiseOps_ex_dispatch_0_cached mutable : !iree.ref<!hal.executable> | |
| vm.func @hloElementwiseOps_ex_dispatch_0(%arg0: !iree.ref<!hal.device>) -> !iree.ref<!hal.executable> { | |
| %hloElementwiseOps_ex_dispatch_0_cached = vm.global.load.ref @hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable> | |
| %rnz = vm.cmp.nz.ref %hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable> | |
| vm.cond_br %rnz, ^bb1(%hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable>), ^bb2 | |
| ^bb1(%0: !iree.ref<!hal.executable>): // pred: ^bb0 | |
| vm.return %0 : !iree.ref<!hal.executable> | |
| ^bb2: // pred: ^bb0 | |
| %c1230128453 = vm.const.i32 1230128453 : i32 | |
| %c1397773893 = vm.const.i32 1397773893 : i32 | |
| %1 = vm.call.variadic @hal.ex.match_supported_executable_format(%arg0, [%c1230128453, %c1397773893]) : (!iree.ref<!hal.device>, i32...) -> i32 | |
| vm.br ^bb3(%1 : i32) | |
| ^bb3(%2: i32): // pred: ^bb2 | |
| %c1230128453_0 = vm.const.i32 1230128453 : i32 | |
| %eq = vm.cmp.eq.i32 %2, %c1230128453_0 : i32 | |
| vm.cond_br %eq, ^bb4(%2 : i32), ^bb5(%2 : i32) | |
| ^bb4(%3: i32): // pred: ^bb3 | |
| %hloElementwiseOps_ex_dispatch_0_data_1230128453 = vm.const.ref.rodata @hloElementwiseOps_ex_dispatch_0_data_1230128453 : !iree.byte_buffer_ref | |
| %ref = vm.call @hal.ex.cache_executable(%arg0, %3, %hloElementwiseOps_ex_dispatch_0_data_1230128453) : (!iree.ref<!hal.device>, i32, !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> | |
| vm.br ^bb7(%ref : !iree.ref<!hal.executable>) | |
| ^bb5(%4: i32): // pred: ^bb3 | |
| %c1397773893_1 = vm.const.i32 1397773893 : i32 | |
| %eq_2 = vm.cmp.eq.i32 %4, %c1397773893_1 : i32 | |
| vm.cond_br %eq_2, ^bb6(%4 : i32), ^bb8 | |
| ^bb6(%5: i32): // pred: ^bb5 | |
| %hloElementwiseOps_ex_dispatch_0_data_1397773893 = vm.const.ref.rodata @hloElementwiseOps_ex_dispatch_0_data_1397773893 : !iree.byte_buffer_ref | |
| %ref_3 = vm.call @hal.ex.cache_executable(%arg0, %5, %hloElementwiseOps_ex_dispatch_0_data_1397773893) : (!iree.ref<!hal.device>, i32, !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> | |
| vm.br ^bb7(%ref_3 : !iree.ref<!hal.executable>) | |
| ^bb7(%6: !iree.ref<!hal.executable>): // 2 preds: ^bb4, ^bb6 | |
| vm.global.store.ref @hloElementwiseOps_ex_dispatch_0_cached, %6 : !iree.ref<!hal.executable> | |
| vm.return %6 : !iree.ref<!hal.executable> | |
| ^bb8: // pred: ^bb5 | |
| %null = vm.const.ref.zero : !iree.ref<!hal.executable> | |
| vm.return %null : !iree.ref<!hal.executable> | |
| } | |
| vm.func @hloElementwiseOps(%arg0: !iree.ref<!hal.buffer>) -> !iree.ref<!hal.buffer> attributes {iree.reflection = {f = "I6!B3!d4R6!B3!d4", fv = "1"}} { | |
| %c1 = vm.const.i32 1 : i32 | |
| %c4 = vm.const.i32 4 : i32 | |
| %ref = vm.call @hal.ex.shared_device() : () -> !iree.ref<!hal.device> | |
| %ref_0 = vm.call @hal.device.allocator(%ref) : (!iree.ref<!hal.device>) -> !iree.ref<!hal.allocator> | |
| %c50 = vm.const.i32 50 : i32 | |
| %c15 = vm.const.i32 15 : i32 | |
| %c4_1 = vm.const.i32 4 : i32 | |
| %ref_2 = vm.call.variadic @hal.allocator.allocate.shaped(%ref_0, %c50, %c15, [%c4], %c4_1) : (!iree.ref<!hal.allocator>, i32, i32, i32..., i32) -> !iree.ref<!hal.buffer> | |
| vm.call @hal.ex.defer_release(%ref_2) : (!iree.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) : (!iree.ref<!hal.device>, i32, i32) -> !iree.ref<!hal.command_buffer> | |
| vm.call @hal.command_buffer.begin(%ref_4) : (!iree.ref<!hal.command_buffer>) -> () | |
| %ref_5 = vm.call @hloElementwiseOps_ex_dispatch_0(%ref) : (!iree.ref<!hal.device>) -> !iree.ref<!hal.executable> | |
| %zero = vm.const.i32.zero : i32 | |
| %c4_6 = vm.const.i32 4 : i32 | |
| vm.call.variadic @hal.ex.push_binding(%ref_4, %zero, %arg0, [%c4], %c4_6) : (!iree.ref<!hal.command_buffer>, i32, !iree.ref<!hal.buffer>, i32..., i32) | |
| vm.call @hal.ex.defer_release(%arg0) : (!iree.ref<!hal.buffer>) -> () | |
| %c1_7 = vm.const.i32 1 : i32 | |
| %c4_8 = vm.const.i32 4 : i32 | |
| vm.call.variadic @hal.ex.push_binding(%ref_4, %c1_7, %ref_2, [%c4], %c4_8) : (!iree.ref<!hal.command_buffer>, i32, !iree.ref<!hal.buffer>, i32..., i32) | |
| vm.call @hal.ex.defer_release(%ref_2) : (!iree.ref<!hal.buffer>) -> () | |
| %zero_9 = vm.const.i32.zero : i32 | |
| vm.call @hal.command_buffer.dispatch(%ref_4, %ref_5, %zero_9, %c1, %c1, %c1) : (!iree.ref<!hal.command_buffer>, !iree.ref<!hal.executable>, i32, i32, i32, i32) -> () | |
| %c16 = vm.const.i32 16 : i32 | |
| %c1_10 = vm.const.i32 1 : i32 | |
| %c8 = vm.const.i32 8 : i32 | |
| %c4_11 = vm.const.i32 4 : i32 | |
| vm.call.variadic @hal.command_buffer.execution_barrier(%ref_4, %c16, %c1_10, [%c8], []) : (!iree.ref<!hal.command_buffer>, i32, i32, i32..., i32...) | |
| vm.call @hal.command_buffer.end(%ref_4) : (!iree.ref<!hal.command_buffer>) -> () | |
| vm.call @hal.ex.submit_and_wait(%ref, %ref_4) : (!iree.ref<!hal.device>, !iree.ref<!hal.command_buffer>) -> () | |
| vm.return %ref_2 : !iree.ref<!hal.buffer> | |
| } | |
| vm.export @hloElementwiseOps | |
| vm.import @hal.ex.shared_device() -> !iree.ref<!hal.device> attributes {nosideeffects} | |
| vm.import @hal.ex.match_supported_executable_format(%device : !iree.ref<!hal.device>, %available_formats : i32...) -> i32 attributes {nosideeffects} | |
| vm.import @hal.ex.cache_executable(%device : !iree.ref<!hal.device>, %executable_format : i32, %executable_data : !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> attributes {nosideeffects} | |
| vm.import @hal.ex.push_binding(%command_buffer : !iree.ref<!hal.command_buffer>, %ordinal : i32, %buffer : !iree.ref<!hal.buffer>, %shape : i32..., %element_size : i32) | |
| vm.import @hal.ex.executable_descriptor_set_layout(%executable : !iree.ref<!hal.executable>, %set : i32) -> !iree.ref<!hal.descriptor_set_layout> | |
| vm.import @hal.ex.defer_release(%operand : !iree.opaque_ref) | |
| vm.import @hal.ex.submit_and_wait(%device : !iree.ref<!hal.device>, %command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.allocator.compute_size(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32) -> i32 attributes {nosideeffects} | |
| vm.import @hal.allocator.allocate(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.allocator.allocate.const(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32, %value : !iree.byte_buffer_ref) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.allocator.allocate.shaped(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.buffer.subspan(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.buffer.fill(%target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) | |
| vm.import @hal.buffer.read_data(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.mutable_byte_buffer_ref, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer.write_data(%target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %source_buffer : !iree.byte_buffer_ref, %source_offset : i32, %length : i32) | |
| vm.import @hal.buffer.copy_data(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer.load(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 | |
| vm.import @hal.buffer.store(%value : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer_view.compute_offset(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %element_size : i32) -> i32 | |
| vm.import @hal.buffer_view.compute_length(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %element_size : i32) -> i32 | |
| vm.import @hal.buffer_view.compute_range(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %lengths : i32..., %element_size : i32) -> (i32, i32) | |
| vm.import @hal.buffer_view.slice(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %lengths : i32..., %element_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.command_buffer.create(%device : !iree.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !iree.ref<!hal.command_buffer> | |
| vm.import @hal.command_buffer.begin(%command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.command_buffer.end(%command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.command_buffer.execution_barrier(%command_buffer : !iree.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) | |
| vm.import @hal.command_buffer.fill_buffer(%command_buffer : !iree.ref<!hal.command_buffer>, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) | |
| vm.import @hal.command_buffer.copy_buffer(%command_buffer : !iree.ref<!hal.command_buffer>, %source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.command_buffer.bind_descriptor_set(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %set : i32, %descriptor_set : !iree.ref<!hal.descriptor_set>, %dynamic_offsets : i32...) | |
| vm.import @hal.command_buffer.dispatch(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) | |
| vm.import @hal.command_buffer.dispatch.indirect(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %entry_point : i32, %workgroups_buffer : !iree.ref<!hal.buffer>, %workgroups_offset : i32) | |
| vm.import @hal.descriptor_set.allocate(%device : !iree.ref<!hal.device>, %set_layout : !iree.ref<!hal.descriptor_set_layout>) -> !iree.ref<!hal.descriptor_set> | |
| vm.import @hal.descriptor_set.update(%device : !iree.ref<!hal.device>, %set : !iree.ref<!hal.descriptor_set>, %binding : i32, %buffer : !iree.ref<!hal.buffer>, %offset : i32, %length : i32, %access : i32) | |
| vm.import @hal.device.allocator(%device : !iree.ref<!hal.device>) -> !iree.ref<!hal.allocator> attributes {nosideeffects} | |
| } | |
| *** IR Dump After CSE *** | |
| module @hal_usage { | |
| vm.module @hal_usage { | |
| vm.rodata @hloElementwiseOps_ex_dispatch_0_data_1230128453 dense<[16, 0, 0, 0, 69, 77, 79, 68, 8, 0, 14, 0, 8, 0, 4, 0, 8, 0, 0, 0, 20, 0, 0, 0, 52, 0, 0, 0, 0, 0, 10, 0, 16, 0, 12, 0, 8, 0, 4, 0, 10, 0, 0, 0, 12, 0, 0, 0, 16, 0, 0, 0, 16, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 28, 0, 0, 0, 6, 0, 0, 0, 109, 111, 100, 117, 108, 101, 0, 0, 12, 0, 16, 0, 12, 0, 8, 0, 0, 0, 4, 0, 12, 0, 0, 0, -52, 0, 0, 0, 48, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 126, -1, -1, -1, 4, 0, 0, 0, 2, 0, 0, 0, 68, 0, 0, 0, 4, 0, 0, 0, -52, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -112, -1, -1, -1, 8, 0, 0, 0, 12, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, -20, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -66, -1, -1, -1, 32, 0, 0, 0, 8, 0, 12, 0, 11, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 1, -52, -1, -1, -1, 8, 0, 0, 0, 20, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, 8, 0, 10, 0, 9, 0, 4, 0, 8, 0, 0, 0, 12, 0, 0, 0, 0, 1, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 32, 0, 0, 0, 8, 0, 12, 0, 8, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 7, 0, 0, 0, 112, 0, 0, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 2, 0, 113, 0, 0, 0, 0, 2, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 3, 0, 115, 2, 0, 0, 0, 3, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 4, 0, 119, 3, 0, 0, 0, 4, 0, 0, 3, 1, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 5, 0, 0, 3, 1, 1, 0, 0, 0, 1, 4, 0, 0, 0, 0, 0, 0, 0, 6, 0, 54, 4, 0, 5, 0, 1, 0, 5, 0, 6, 0, 4, 0]> : vector<436xi8> | |
| vm.rodata @hloElementwiseOps_ex_dispatch_0_data_1397773893 dense<[8, 0, 0, 0, 83, 80, 86, 69, -92, -1, -1, -1, 104, 4, 0, 0, 112, 0, 0, 0, 4, 0, 0, 0, -18, -1, -1, -1, 4, 0, 0, 0, 1, 0, 0, 0, 12, 0, 0, 0, 0, 0, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 52, 0, 0, 0, 16, 0, 0, 0, 12, 0, 20, 0, 4, 0, 8, 0, 12, 0, 16, 0, 12, 0, 0, 0, 1, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, 12, 0, 16, 0, 0, 0, 4, 0, 8, 0, 12, 0, 12, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, -4, 0, 0, 0, 3, 2, 35, 7, 0, 0, 1, 0, 22, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 17, 0, 2, 0, 1, 0, 0, 0, 10, 0, 11, 0, 83, 80, 86, 95, 75, 72, 82, 95, 115, 116, 111, 114, 97, 103, 101, 95, 98, 117, 102, 102, 101, 114, 95, 115, 116, 111, 114, 97, 103, 101, 95, 99, 108, 97, 115, 115, 0, 0, 0, 0, 14, 0, 3, 0, 0, 0, 0, 0, 1, 0, 0, 0, 15, 0, 13, 0, 5, 0, 0, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 4, 0, 0, 0, 16, 0, 6, 0, 14, 0, 0, 0, 17, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 5, 0, 7, 0, 4, 0, 0, 0, 103, 108, 111, 98, 97, 108, 73, 110, 118, 111, 99, 97, 116, 105, 111, 110, 73, 68, 0, 0, 5, 0, 12, 0, 10, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 48, 0, 0, 5, 0, 12, 0, 11, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 49, 0, 0, 5, 0, 11, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 71, 0, 4, 0, 4, 0, 0, 0, 11, 0, 0, 0, 28, 0, 0, 0, 71, 0, 4, 0, 7, 0, 0, 0, 6, 0, 0, 0, 4, 0, 0, 0, 72, 0, 5, 0, 6, 0, 0, 0, 0, 0, 0, 0, 35, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 33, 0, 0, 0, 1, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 21, 0, 4, 0, 3, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 23, 0, 4, 0, 2, 0, 0, 0, 3, 0, 0, 0, 3, 0, 0, 0, 32, 0, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 2, 0, 0, 0, 59, 0, 4, 0, 1, 0, 0, 0, 4, 0, 0, 0, 1, 0, 0, 0, 22, 0, 3, 0, 8, 0, 0, 0, 32, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 9, 0, 0, 0, 4, 0, 0, 0, 28, 0, 4, 0, 7, 0, 0, 0, 8, 0, 0, 0, 9, 0, 0, 0, 30, 0, 3, 0, 6, 0, 0, 0, 7, 0, 0, 0, 32, 0, 4, 0, 5, 0, 0, 0, 12, 0, 0, 0, 6, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 10, 0, 0, 0, 12, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 11, 0, 0, 0, 12, 0, 0, 0, 19, 0, 2, 0, 13, 0, 0, 0, 33, 0, 3, 0, 12, 0, 0, 0, 13, 0, 0, 0, 20, 0, 2, 0, 16, 0, 0, 0, 41, 0, 3, 0, 16, 0, 0, 0, 17, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 18, 0, 0, 0, 0, 0, 0, 0, 32, 0, 4, 0, 26, 0, 0, 0, 12, 0, 0, 0, 8, 0, 0, 0, 54, 0, 5, 0, 13, 0, 0, 0, 14, 0, 0, 0, 0, 0, 0, 0, 12, 0, 0, 0, -8, 0, 2, 0, 15, 0, 0, 0, 61, 0, 4, 0, 2, 0, 0, 0, 19, 0, 0, 0, 4, 0, 0, 0, 81, 0, 5, 0, 3, 0, 0, 0, 20, 0, 0, 0, 19, 0, 0, 0, 0, 0, 0, 0, -79, 0, 5, 0, 16, 0, 0, 0, 21, 0, 0, 0, 20, 0, 0, 0, 9, 0, 0, 0, -89, 0, 5, 0, 16, 0, 0, 0, 22, 0, 0, 0, 21, 0, 0, 0, 17, 0, 0, 0, -9, 0, 3, 0, 25, 0, 0, 0, 0, 0, 0, 0, -6, 0, 4, 0, 22, 0, 0, 0, 24, 0, 0, 0, 25, 0, 0, 0, -8, 0, 2, 0, 24, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 27, 0, 0, 0, 10, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 61, 0, 4, 0, 8, 0, 0, 0, 28, 0, 0, 0, 27, 0, 0, 0, -127, 0, 5, 0, 8, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, 28, 0, 0, 0, -125, 0, 5, 0, 8, 0, 0, 0, 30, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, -123, 0, 5, 0, 8, 0, 0, 0, 31, 0, 0, 0, 30, 0, 0, 0, 28, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 32, 0, 0, 0, 11, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 62, 0, 3, 0, 32, 0, 0, 0, 31, 0, 0, 0, -7, 0, 2, 0, 25, 0, 0, 0, -8, 0, 2, 0, 25, 0, 0, 0, -3, 0, 1, 0, 56, 0, 1, 0, 1, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0]> : vector<1188xi8> | |
| vm.global.ref @hloElementwiseOps_ex_dispatch_0_cached mutable : !iree.ref<!hal.executable> | |
| vm.func @hloElementwiseOps_ex_dispatch_0(%arg0: !iree.ref<!hal.device>) -> !iree.ref<!hal.executable> { | |
| %hloElementwiseOps_ex_dispatch_0_cached = vm.global.load.ref @hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable> | |
| %rnz = vm.cmp.nz.ref %hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable> | |
| vm.cond_br %rnz, ^bb1(%hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable>), ^bb2 | |
| ^bb1(%0: !iree.ref<!hal.executable>): // pred: ^bb0 | |
| vm.return %0 : !iree.ref<!hal.executable> | |
| ^bb2: // pred: ^bb0 | |
| %c1230128453 = vm.const.i32 1230128453 : i32 | |
| %c1397773893 = vm.const.i32 1397773893 : i32 | |
| %1 = vm.call.variadic @hal.ex.match_supported_executable_format(%arg0, [%c1230128453, %c1397773893]) : (!iree.ref<!hal.device>, i32...) -> i32 | |
| vm.br ^bb3(%1 : i32) | |
| ^bb3(%2: i32): // pred: ^bb2 | |
| %eq = vm.cmp.eq.i32 %2, %c1230128453 : i32 | |
| vm.cond_br %eq, ^bb4(%2 : i32), ^bb5(%2 : i32) | |
| ^bb4(%3: i32): // pred: ^bb3 | |
| %hloElementwiseOps_ex_dispatch_0_data_1230128453 = vm.const.ref.rodata @hloElementwiseOps_ex_dispatch_0_data_1230128453 : !iree.byte_buffer_ref | |
| %ref = vm.call @hal.ex.cache_executable(%arg0, %3, %hloElementwiseOps_ex_dispatch_0_data_1230128453) : (!iree.ref<!hal.device>, i32, !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> | |
| vm.br ^bb7(%ref : !iree.ref<!hal.executable>) | |
| ^bb5(%4: i32): // pred: ^bb3 | |
| %eq_0 = vm.cmp.eq.i32 %4, %c1397773893 : i32 | |
| vm.cond_br %eq_0, ^bb6(%4 : i32), ^bb8 | |
| ^bb6(%5: i32): // pred: ^bb5 | |
| %hloElementwiseOps_ex_dispatch_0_data_1397773893 = vm.const.ref.rodata @hloElementwiseOps_ex_dispatch_0_data_1397773893 : !iree.byte_buffer_ref | |
| %ref_1 = vm.call @hal.ex.cache_executable(%arg0, %5, %hloElementwiseOps_ex_dispatch_0_data_1397773893) : (!iree.ref<!hal.device>, i32, !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> | |
| vm.br ^bb7(%ref_1 : !iree.ref<!hal.executable>) | |
| ^bb7(%6: !iree.ref<!hal.executable>): // 2 preds: ^bb4, ^bb6 | |
| vm.global.store.ref @hloElementwiseOps_ex_dispatch_0_cached, %6 : !iree.ref<!hal.executable> | |
| vm.return %6 : !iree.ref<!hal.executable> | |
| ^bb8: // pred: ^bb5 | |
| %null = vm.const.ref.zero : !iree.ref<!hal.executable> | |
| vm.return %null : !iree.ref<!hal.executable> | |
| } | |
| vm.func @hloElementwiseOps(%arg0: !iree.ref<!hal.buffer>) -> !iree.ref<!hal.buffer> attributes {iree.reflection = {f = "I6!B3!d4R6!B3!d4", fv = "1"}} { | |
| %c1 = vm.const.i32 1 : i32 | |
| %c4 = vm.const.i32 4 : i32 | |
| %ref = vm.call @hal.ex.shared_device() : () -> !iree.ref<!hal.device> | |
| %ref_0 = vm.call @hal.device.allocator(%ref) : (!iree.ref<!hal.device>) -> !iree.ref<!hal.allocator> | |
| %c50 = vm.const.i32 50 : i32 | |
| %c15 = vm.const.i32 15 : i32 | |
| %ref_1 = vm.call.variadic @hal.allocator.allocate.shaped(%ref_0, %c50, %c15, [%c4], %c4) : (!iree.ref<!hal.allocator>, i32, i32, i32..., i32) -> !iree.ref<!hal.buffer> | |
| vm.call @hal.ex.defer_release(%ref_1) : (!iree.ref<!hal.buffer>) -> () | |
| %c3 = vm.const.i32 3 : i32 | |
| %ref_2 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!iree.ref<!hal.device>, i32, i32) -> !iree.ref<!hal.command_buffer> | |
| vm.call @hal.command_buffer.begin(%ref_2) : (!iree.ref<!hal.command_buffer>) -> () | |
| %ref_3 = vm.call @hloElementwiseOps_ex_dispatch_0(%ref) : (!iree.ref<!hal.device>) -> !iree.ref<!hal.executable> | |
| %zero = vm.const.i32.zero : i32 | |
| vm.call.variadic @hal.ex.push_binding(%ref_2, %zero, %arg0, [%c4], %c4) : (!iree.ref<!hal.command_buffer>, i32, !iree.ref<!hal.buffer>, i32..., i32) | |
| vm.call @hal.ex.defer_release(%arg0) : (!iree.ref<!hal.buffer>) -> () | |
| vm.call.variadic @hal.ex.push_binding(%ref_2, %c1, %ref_1, [%c4], %c4) : (!iree.ref<!hal.command_buffer>, i32, !iree.ref<!hal.buffer>, i32..., i32) | |
| vm.call @hal.ex.defer_release(%ref_1) : (!iree.ref<!hal.buffer>) -> () | |
| vm.call @hal.command_buffer.dispatch(%ref_2, %ref_3, %zero, %c1, %c1, %c1) : (!iree.ref<!hal.command_buffer>, !iree.ref<!hal.executable>, i32, i32, i32, i32) -> () | |
| %c16 = vm.const.i32 16 : i32 | |
| %c8 = vm.const.i32 8 : i32 | |
| vm.call.variadic @hal.command_buffer.execution_barrier(%ref_2, %c16, %c1, [%c8], []) : (!iree.ref<!hal.command_buffer>, i32, i32, i32..., i32...) | |
| vm.call @hal.command_buffer.end(%ref_2) : (!iree.ref<!hal.command_buffer>) -> () | |
| vm.call @hal.ex.submit_and_wait(%ref, %ref_2) : (!iree.ref<!hal.device>, !iree.ref<!hal.command_buffer>) -> () | |
| vm.return %ref_1 : !iree.ref<!hal.buffer> | |
| } | |
| vm.export @hloElementwiseOps | |
| vm.import @hal.ex.shared_device() -> !iree.ref<!hal.device> attributes {nosideeffects} | |
| vm.import @hal.ex.match_supported_executable_format(%device : !iree.ref<!hal.device>, %available_formats : i32...) -> i32 attributes {nosideeffects} | |
| vm.import @hal.ex.cache_executable(%device : !iree.ref<!hal.device>, %executable_format : i32, %executable_data : !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> attributes {nosideeffects} | |
| vm.import @hal.ex.push_binding(%command_buffer : !iree.ref<!hal.command_buffer>, %ordinal : i32, %buffer : !iree.ref<!hal.buffer>, %shape : i32..., %element_size : i32) | |
| vm.import @hal.ex.executable_descriptor_set_layout(%executable : !iree.ref<!hal.executable>, %set : i32) -> !iree.ref<!hal.descriptor_set_layout> | |
| vm.import @hal.ex.defer_release(%operand : !iree.opaque_ref) | |
| vm.import @hal.ex.submit_and_wait(%device : !iree.ref<!hal.device>, %command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.allocator.compute_size(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32) -> i32 attributes {nosideeffects} | |
| vm.import @hal.allocator.allocate(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.allocator.allocate.const(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32, %value : !iree.byte_buffer_ref) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.allocator.allocate.shaped(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.buffer.subspan(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.buffer.fill(%target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) | |
| vm.import @hal.buffer.read_data(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.mutable_byte_buffer_ref, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer.write_data(%target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %source_buffer : !iree.byte_buffer_ref, %source_offset : i32, %length : i32) | |
| vm.import @hal.buffer.copy_data(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer.load(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 | |
| vm.import @hal.buffer.store(%value : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer_view.compute_offset(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %element_size : i32) -> i32 | |
| vm.import @hal.buffer_view.compute_length(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %element_size : i32) -> i32 | |
| vm.import @hal.buffer_view.compute_range(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %lengths : i32..., %element_size : i32) -> (i32, i32) | |
| vm.import @hal.buffer_view.slice(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %lengths : i32..., %element_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.command_buffer.create(%device : !iree.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !iree.ref<!hal.command_buffer> | |
| vm.import @hal.command_buffer.begin(%command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.command_buffer.end(%command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.command_buffer.execution_barrier(%command_buffer : !iree.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) | |
| vm.import @hal.command_buffer.fill_buffer(%command_buffer : !iree.ref<!hal.command_buffer>, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) | |
| vm.import @hal.command_buffer.copy_buffer(%command_buffer : !iree.ref<!hal.command_buffer>, %source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.command_buffer.bind_descriptor_set(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %set : i32, %descriptor_set : !iree.ref<!hal.descriptor_set>, %dynamic_offsets : i32...) | |
| vm.import @hal.command_buffer.dispatch(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) | |
| vm.import @hal.command_buffer.dispatch.indirect(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %entry_point : i32, %workgroups_buffer : !iree.ref<!hal.buffer>, %workgroups_offset : i32) | |
| vm.import @hal.descriptor_set.allocate(%device : !iree.ref<!hal.device>, %set_layout : !iree.ref<!hal.descriptor_set_layout>) -> !iree.ref<!hal.descriptor_set> | |
| vm.import @hal.descriptor_set.update(%device : !iree.ref<!hal.device>, %set : !iree.ref<!hal.descriptor_set>, %binding : i32, %buffer : !iree.ref<!hal.buffer>, %offset : i32, %length : i32, %access : i32) | |
| vm.import @hal.device.allocator(%device : !iree.ref<!hal.device>) -> !iree.ref<!hal.allocator> attributes {nosideeffects} | |
| } | |
| } | |
| *** IR Dump After iree_compiler::IREE::DropCompilerHintsPass *** | |
| module @hal_usage { | |
| vm.module @hal_usage { | |
| vm.rodata @hloElementwiseOps_ex_dispatch_0_data_1230128453 dense<[16, 0, 0, 0, 69, 77, 79, 68, 8, 0, 14, 0, 8, 0, 4, 0, 8, 0, 0, 0, 20, 0, 0, 0, 52, 0, 0, 0, 0, 0, 10, 0, 16, 0, 12, 0, 8, 0, 4, 0, 10, 0, 0, 0, 12, 0, 0, 0, 16, 0, 0, 0, 16, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 28, 0, 0, 0, 6, 0, 0, 0, 109, 111, 100, 117, 108, 101, 0, 0, 12, 0, 16, 0, 12, 0, 8, 0, 0, 0, 4, 0, 12, 0, 0, 0, -52, 0, 0, 0, 48, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 126, -1, -1, -1, 4, 0, 0, 0, 2, 0, 0, 0, 68, 0, 0, 0, 4, 0, 0, 0, -52, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -112, -1, -1, -1, 8, 0, 0, 0, 12, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, -20, -1, -1, -1, 8, 0, 0, 0, 0, 0, 0, 1, -66, -1, -1, -1, 32, 0, 0, 0, 8, 0, 12, 0, 11, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 1, -52, -1, -1, -1, 8, 0, 0, 0, 20, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, 8, 0, 10, 0, 9, 0, 4, 0, 8, 0, 0, 0, 12, 0, 0, 0, 0, 1, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 32, 0, 0, 0, 8, 0, 12, 0, 8, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 7, 0, 0, 0, 112, 0, 0, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 2, 0, 113, 0, 0, 0, 0, 2, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 3, 0, 115, 2, 0, 0, 0, 3, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 4, 0, 119, 3, 0, 0, 0, 4, 0, 0, 3, 1, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 5, 0, 0, 3, 1, 1, 0, 0, 0, 1, 4, 0, 0, 0, 0, 0, 0, 0, 6, 0, 54, 4, 0, 5, 0, 1, 0, 5, 0, 6, 0, 4, 0]> : vector<436xi8> | |
| vm.rodata @hloElementwiseOps_ex_dispatch_0_data_1397773893 dense<[8, 0, 0, 0, 83, 80, 86, 69, -92, -1, -1, -1, 104, 4, 0, 0, 112, 0, 0, 0, 4, 0, 0, 0, -18, -1, -1, -1, 4, 0, 0, 0, 1, 0, 0, 0, 12, 0, 0, 0, 0, 0, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 52, 0, 0, 0, 16, 0, 0, 0, 12, 0, 20, 0, 4, 0, 8, 0, 12, 0, 16, 0, 12, 0, 0, 0, 1, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, 12, 0, 16, 0, 0, 0, 4, 0, 8, 0, 12, 0, 12, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, -4, 0, 0, 0, 3, 2, 35, 7, 0, 0, 1, 0, 22, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 17, 0, 2, 0, 1, 0, 0, 0, 10, 0, 11, 0, 83, 80, 86, 95, 75, 72, 82, 95, 115, 116, 111, 114, 97, 103, 101, 95, 98, 117, 102, 102, 101, 114, 95, 115, 116, 111, 114, 97, 103, 101, 95, 99, 108, 97, 115, 115, 0, 0, 0, 0, 14, 0, 3, 0, 0, 0, 0, 0, 1, 0, 0, 0, 15, 0, 13, 0, 5, 0, 0, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 4, 0, 0, 0, 16, 0, 6, 0, 14, 0, 0, 0, 17, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 5, 0, 7, 0, 4, 0, 0, 0, 103, 108, 111, 98, 97, 108, 73, 110, 118, 111, 99, 97, 116, 105, 111, 110, 73, 68, 0, 0, 5, 0, 12, 0, 10, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 48, 0, 0, 5, 0, 12, 0, 11, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 49, 0, 0, 5, 0, 11, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 71, 0, 4, 0, 4, 0, 0, 0, 11, 0, 0, 0, 28, 0, 0, 0, 71, 0, 4, 0, 7, 0, 0, 0, 6, 0, 0, 0, 4, 0, 0, 0, 72, 0, 5, 0, 6, 0, 0, 0, 0, 0, 0, 0, 35, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 33, 0, 0, 0, 1, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 21, 0, 4, 0, 3, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 23, 0, 4, 0, 2, 0, 0, 0, 3, 0, 0, 0, 3, 0, 0, 0, 32, 0, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 2, 0, 0, 0, 59, 0, 4, 0, 1, 0, 0, 0, 4, 0, 0, 0, 1, 0, 0, 0, 22, 0, 3, 0, 8, 0, 0, 0, 32, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 9, 0, 0, 0, 4, 0, 0, 0, 28, 0, 4, 0, 7, 0, 0, 0, 8, 0, 0, 0, 9, 0, 0, 0, 30, 0, 3, 0, 6, 0, 0, 0, 7, 0, 0, 0, 32, 0, 4, 0, 5, 0, 0, 0, 12, 0, 0, 0, 6, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 10, 0, 0, 0, 12, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 11, 0, 0, 0, 12, 0, 0, 0, 19, 0, 2, 0, 13, 0, 0, 0, 33, 0, 3, 0, 12, 0, 0, 0, 13, 0, 0, 0, 20, 0, 2, 0, 16, 0, 0, 0, 41, 0, 3, 0, 16, 0, 0, 0, 17, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 18, 0, 0, 0, 0, 0, 0, 0, 32, 0, 4, 0, 26, 0, 0, 0, 12, 0, 0, 0, 8, 0, 0, 0, 54, 0, 5, 0, 13, 0, 0, 0, 14, 0, 0, 0, 0, 0, 0, 0, 12, 0, 0, 0, -8, 0, 2, 0, 15, 0, 0, 0, 61, 0, 4, 0, 2, 0, 0, 0, 19, 0, 0, 0, 4, 0, 0, 0, 81, 0, 5, 0, 3, 0, 0, 0, 20, 0, 0, 0, 19, 0, 0, 0, 0, 0, 0, 0, -79, 0, 5, 0, 16, 0, 0, 0, 21, 0, 0, 0, 20, 0, 0, 0, 9, 0, 0, 0, -89, 0, 5, 0, 16, 0, 0, 0, 22, 0, 0, 0, 21, 0, 0, 0, 17, 0, 0, 0, -9, 0, 3, 0, 25, 0, 0, 0, 0, 0, 0, 0, -6, 0, 4, 0, 22, 0, 0, 0, 24, 0, 0, 0, 25, 0, 0, 0, -8, 0, 2, 0, 24, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 27, 0, 0, 0, 10, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 61, 0, 4, 0, 8, 0, 0, 0, 28, 0, 0, 0, 27, 0, 0, 0, -127, 0, 5, 0, 8, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, 28, 0, 0, 0, -125, 0, 5, 0, 8, 0, 0, 0, 30, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, -123, 0, 5, 0, 8, 0, 0, 0, 31, 0, 0, 0, 30, 0, 0, 0, 28, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 32, 0, 0, 0, 11, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 62, 0, 3, 0, 32, 0, 0, 0, 31, 0, 0, 0, -7, 0, 2, 0, 25, 0, 0, 0, -8, 0, 2, 0, 25, 0, 0, 0, -3, 0, 1, 0, 56, 0, 1, 0, 1, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0]> : vector<1188xi8> | |
| vm.global.ref @hloElementwiseOps_ex_dispatch_0_cached mutable : !iree.ref<!hal.executable> | |
| vm.func @hloElementwiseOps_ex_dispatch_0(%arg0: !iree.ref<!hal.device>) -> !iree.ref<!hal.executable> { | |
| %hloElementwiseOps_ex_dispatch_0_cached = vm.global.load.ref @hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable> | |
| %rnz = vm.cmp.nz.ref %hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable> | |
| vm.cond_br %rnz, ^bb1(%hloElementwiseOps_ex_dispatch_0_cached : !iree.ref<!hal.executable>), ^bb2 | |
| ^bb1(%0: !iree.ref<!hal.executable>): // pred: ^bb0 | |
| vm.return %0 : !iree.ref<!hal.executable> | |
| ^bb2: // pred: ^bb0 | |
| %c1230128453 = vm.const.i32 1230128453 : i32 | |
| %c1397773893 = vm.const.i32 1397773893 : i32 | |
| %1 = vm.call.variadic @hal.ex.match_supported_executable_format(%arg0, [%c1230128453, %c1397773893]) : (!iree.ref<!hal.device>, i32...) -> i32 | |
| vm.br ^bb3(%1 : i32) | |
| ^bb3(%2: i32): // pred: ^bb2 | |
| %eq = vm.cmp.eq.i32 %2, %c1230128453 : i32 | |
| vm.cond_br %eq, ^bb4(%2 : i32), ^bb5(%2 : i32) | |
| ^bb4(%3: i32): // pred: ^bb3 | |
| %hloElementwiseOps_ex_dispatch_0_data_1230128453 = vm.const.ref.rodata @hloElementwiseOps_ex_dispatch_0_data_1230128453 : !iree.byte_buffer_ref | |
| %ref = vm.call @hal.ex.cache_executable(%arg0, %3, %hloElementwiseOps_ex_dispatch_0_data_1230128453) : (!iree.ref<!hal.device>, i32, !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> | |
| vm.br ^bb7(%ref : !iree.ref<!hal.executable>) | |
| ^bb5(%4: i32): // pred: ^bb3 | |
| %eq_0 = vm.cmp.eq.i32 %4, %c1397773893 : i32 | |
| vm.cond_br %eq_0, ^bb6(%4 : i32), ^bb8 | |
| ^bb6(%5: i32): // pred: ^bb5 | |
| %hloElementwiseOps_ex_dispatch_0_data_1397773893 = vm.const.ref.rodata @hloElementwiseOps_ex_dispatch_0_data_1397773893 : !iree.byte_buffer_ref | |
| %ref_1 = vm.call @hal.ex.cache_executable(%arg0, %5, %hloElementwiseOps_ex_dispatch_0_data_1397773893) : (!iree.ref<!hal.device>, i32, !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> | |
| vm.br ^bb7(%ref_1 : !iree.ref<!hal.executable>) | |
| ^bb7(%6: !iree.ref<!hal.executable>): // 2 preds: ^bb4, ^bb6 | |
| vm.global.store.ref @hloElementwiseOps_ex_dispatch_0_cached, %6 : !iree.ref<!hal.executable> | |
| vm.return %6 : !iree.ref<!hal.executable> | |
| ^bb8: // pred: ^bb5 | |
| %null = vm.const.ref.zero : !iree.ref<!hal.executable> | |
| vm.return %null : !iree.ref<!hal.executable> | |
| } | |
| vm.func @hloElementwiseOps(%arg0: !iree.ref<!hal.buffer>) -> !iree.ref<!hal.buffer> attributes {iree.reflection = {f = "I6!B3!d4R6!B3!d4", fv = "1"}} { | |
| %c1 = vm.const.i32 1 : i32 | |
| %c4 = vm.const.i32 4 : i32 | |
| %ref = vm.call @hal.ex.shared_device() : () -> !iree.ref<!hal.device> | |
| %ref_0 = vm.call @hal.device.allocator(%ref) : (!iree.ref<!hal.device>) -> !iree.ref<!hal.allocator> | |
| %c50 = vm.const.i32 50 : i32 | |
| %c15 = vm.const.i32 15 : i32 | |
| %ref_1 = vm.call.variadic @hal.allocator.allocate.shaped(%ref_0, %c50, %c15, [%c4], %c4) : (!iree.ref<!hal.allocator>, i32, i32, i32..., i32) -> !iree.ref<!hal.buffer> | |
| vm.call @hal.ex.defer_release(%ref_1) : (!iree.ref<!hal.buffer>) -> () | |
| %c3 = vm.const.i32 3 : i32 | |
| %ref_2 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!iree.ref<!hal.device>, i32, i32) -> !iree.ref<!hal.command_buffer> | |
| vm.call @hal.command_buffer.begin(%ref_2) : (!iree.ref<!hal.command_buffer>) -> () | |
| %ref_3 = vm.call @hloElementwiseOps_ex_dispatch_0(%ref) : (!iree.ref<!hal.device>) -> !iree.ref<!hal.executable> | |
| %zero = vm.const.i32.zero : i32 | |
| vm.call.variadic @hal.ex.push_binding(%ref_2, %zero, %arg0, [%c4], %c4) : (!iree.ref<!hal.command_buffer>, i32, !iree.ref<!hal.buffer>, i32..., i32) | |
| vm.call @hal.ex.defer_release(%arg0) : (!iree.ref<!hal.buffer>) -> () | |
| vm.call.variadic @hal.ex.push_binding(%ref_2, %c1, %ref_1, [%c4], %c4) : (!iree.ref<!hal.command_buffer>, i32, !iree.ref<!hal.buffer>, i32..., i32) | |
| vm.call @hal.ex.defer_release(%ref_1) : (!iree.ref<!hal.buffer>) -> () | |
| vm.call @hal.command_buffer.dispatch(%ref_2, %ref_3, %zero, %c1, %c1, %c1) : (!iree.ref<!hal.command_buffer>, !iree.ref<!hal.executable>, i32, i32, i32, i32) -> () | |
| %c16 = vm.const.i32 16 : i32 | |
| %c8 = vm.const.i32 8 : i32 | |
| vm.call.variadic @hal.command_buffer.execution_barrier(%ref_2, %c16, %c1, [%c8], []) : (!iree.ref<!hal.command_buffer>, i32, i32, i32..., i32...) | |
| vm.call @hal.command_buffer.end(%ref_2) : (!iree.ref<!hal.command_buffer>) -> () | |
| vm.call @hal.ex.submit_and_wait(%ref, %ref_2) : (!iree.ref<!hal.device>, !iree.ref<!hal.command_buffer>) -> () | |
| vm.return %ref_1 : !iree.ref<!hal.buffer> | |
| } | |
| vm.export @hloElementwiseOps | |
| vm.import @hal.ex.shared_device() -> !iree.ref<!hal.device> attributes {nosideeffects} | |
| vm.import @hal.ex.match_supported_executable_format(%device : !iree.ref<!hal.device>, %available_formats : i32...) -> i32 attributes {nosideeffects} | |
| vm.import @hal.ex.cache_executable(%device : !iree.ref<!hal.device>, %executable_format : i32, %executable_data : !iree.byte_buffer_ref) -> !iree.ref<!hal.executable> attributes {nosideeffects} | |
| vm.import @hal.ex.push_binding(%command_buffer : !iree.ref<!hal.command_buffer>, %ordinal : i32, %buffer : !iree.ref<!hal.buffer>, %shape : i32..., %element_size : i32) | |
| vm.import @hal.ex.executable_descriptor_set_layout(%executable : !iree.ref<!hal.executable>, %set : i32) -> !iree.ref<!hal.descriptor_set_layout> | |
| vm.import @hal.ex.defer_release(%operand : !iree.opaque_ref) | |
| vm.import @hal.ex.submit_and_wait(%device : !iree.ref<!hal.device>, %command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.allocator.compute_size(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32) -> i32 attributes {nosideeffects} | |
| vm.import @hal.allocator.allocate(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.allocator.allocate.const(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32, %value : !iree.byte_buffer_ref) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.allocator.allocate.shaped(%allocator : !iree.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.buffer.subspan(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.buffer.fill(%target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) | |
| vm.import @hal.buffer.read_data(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.mutable_byte_buffer_ref, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer.write_data(%target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %source_buffer : !iree.byte_buffer_ref, %source_offset : i32, %length : i32) | |
| vm.import @hal.buffer.copy_data(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer.load(%source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 | |
| vm.import @hal.buffer.store(%value : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.buffer_view.compute_offset(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %element_size : i32) -> i32 | |
| vm.import @hal.buffer_view.compute_length(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %element_size : i32) -> i32 | |
| vm.import @hal.buffer_view.compute_range(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %lengths : i32..., %element_size : i32) -> (i32, i32) | |
| vm.import @hal.buffer_view.slice(%buffer : !iree.ref<!hal.buffer>, %shape : i32..., %indices : i32..., %lengths : i32..., %element_size : i32) -> !iree.ref<!hal.buffer> | |
| vm.import @hal.command_buffer.create(%device : !iree.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !iree.ref<!hal.command_buffer> | |
| vm.import @hal.command_buffer.begin(%command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.command_buffer.end(%command_buffer : !iree.ref<!hal.command_buffer>) | |
| vm.import @hal.command_buffer.execution_barrier(%command_buffer : !iree.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) | |
| vm.import @hal.command_buffer.fill_buffer(%command_buffer : !iree.ref<!hal.command_buffer>, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) | |
| vm.import @hal.command_buffer.copy_buffer(%command_buffer : !iree.ref<!hal.command_buffer>, %source_buffer : !iree.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !iree.ref<!hal.buffer>, %target_offset : i32, %length : i32) | |
| vm.import @hal.command_buffer.bind_descriptor_set(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %set : i32, %descriptor_set : !iree.ref<!hal.descriptor_set>, %dynamic_offsets : i32...) | |
| vm.import @hal.command_buffer.dispatch(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) | |
| vm.import @hal.command_buffer.dispatch.indirect(%command_buffer : !iree.ref<!hal.command_buffer>, %executable : !iree.ref<!hal.executable>, %entry_point : i32, %workgroups_buffer : !iree.ref<!hal.buffer>, %workgroups_offset : i32) | |
| vm.import @hal.descriptor_set.allocate(%device : !iree.ref<!hal.device>, %set_layout : !iree.ref<!hal.descriptor_set_layout>) -> !iree.ref<!hal.descriptor_set> | |
| vm.import @hal.descriptor_set.update(%device : !iree.ref<!hal.device>, %set : !iree.ref<!hal.descriptor_set>, %binding : i32, %buffer : !iree.ref<!hal.buffer>, %offset : i32, %length : i32, %access : i32) | |
| vm.import @hal.device.allocator(%device : !iree.ref<!hal.device>) -> !iree.ref<!hal.allocator> attributes {nosideeffects} | |
| } | |
| } | |
| { | |
| name: "hal_usage", | |
| types: [ { | |
| full_name: "i32" | |
| }, { | |
| full_name: "!hal.allocator" | |
| }, { | |
| full_name: "!hal.buffer" | |
| }, { | |
| full_name: "!hal.command_buffer" | |
| }, { | |
| full_name: "!hal.device" | |
| }, { | |
| full_name: "!hal.executable" | |
| }, { | |
| full_name: "!iree.byte_buffer" | |
| } ], | |
| imported_functions: [ { | |
| full_name: "hal.ex.shared_device", | |
| signature: { | |
| result_types: [ 4 ] | |
| } | |
| }, { | |
| full_name: "hal.ex.match_supported_executable_format", | |
| signature: { | |
| argument_types: [ 4, 0 ], | |
| result_types: [ 0 ] | |
| } | |
| }, { | |
| full_name: "hal.ex.cache_executable", | |
| signature: { | |
| argument_types: [ 4, 0, 6 ], | |
| result_types: [ 5 ] | |
| } | |
| }, { | |
| full_name: "hal.ex.push_binding", | |
| signature: { | |
| argument_types: [ 3, 0, 2, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.ex.executable_descriptor_set_layout", | |
| signature: { | |
| argument_types: [ 5, 0 ], | |
| result_types: [ 0 ] | |
| } | |
| }, { | |
| full_name: "hal.ex.defer_release", | |
| signature: { | |
| argument_types: [ 0 ] | |
| } | |
| }, { | |
| full_name: "hal.ex.submit_and_wait", | |
| signature: { | |
| argument_types: [ 4, 3 ] | |
| } | |
| }, { | |
| full_name: "hal.allocator.compute_size", | |
| signature: { | |
| argument_types: [ 1, 0, 0, 0, 0 ], | |
| result_types: [ 0 ] | |
| } | |
| }, { | |
| full_name: "hal.allocator.allocate", | |
| signature: { | |
| argument_types: [ 1, 0, 0, 0 ], | |
| result_types: [ 2 ] | |
| } | |
| }, { | |
| full_name: "hal.allocator.allocate.const", | |
| signature: { | |
| argument_types: [ 1, 0, 0, 0, 0, 6 ], | |
| result_types: [ 2 ] | |
| } | |
| }, { | |
| full_name: "hal.allocator.allocate.shaped", | |
| signature: { | |
| argument_types: [ 1, 0, 0, 0, 0 ], | |
| result_types: [ 2 ] | |
| } | |
| }, { | |
| full_name: "hal.buffer.subspan", | |
| signature: { | |
| argument_types: [ 2, 0, 0 ], | |
| result_types: [ 2 ] | |
| } | |
| }, { | |
| full_name: "hal.buffer.fill", | |
| signature: { | |
| argument_types: [ 2, 0, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.buffer.read_data", | |
| signature: { | |
| argument_types: [ 2, 0, 0, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.buffer.write_data", | |
| signature: { | |
| argument_types: [ 2, 0, 6, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.buffer.copy_data", | |
| signature: { | |
| argument_types: [ 2, 0, 2, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.buffer.load", | |
| signature: { | |
| argument_types: [ 2, 0, 0 ], | |
| result_types: [ 0 ] | |
| } | |
| }, { | |
| full_name: "hal.buffer.store", | |
| signature: { | |
| argument_types: [ 0, 2, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.buffer_view.compute_offset", | |
| signature: { | |
| argument_types: [ 2, 0, 0, 0 ], | |
| result_types: [ 0 ] | |
| } | |
| }, { | |
| full_name: "hal.buffer_view.compute_length", | |
| signature: { | |
| argument_types: [ 2, 0, 0 ], | |
| result_types: [ 0 ] | |
| } | |
| }, { | |
| full_name: "hal.buffer_view.compute_range", | |
| signature: { | |
| argument_types: [ 2, 0, 0, 0, 0 ], | |
| result_types: [ 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.buffer_view.slice", | |
| signature: { | |
| argument_types: [ 2, 0, 0, 0, 0 ], | |
| result_types: [ 2 ] | |
| } | |
| }, { | |
| full_name: "hal.command_buffer.create", | |
| signature: { | |
| argument_types: [ 4, 0, 0 ], | |
| result_types: [ 3 ] | |
| } | |
| }, { | |
| full_name: "hal.command_buffer.begin", | |
| signature: { | |
| argument_types: [ 3 ] | |
| } | |
| }, { | |
| full_name: "hal.command_buffer.end", | |
| signature: { | |
| argument_types: [ 3 ] | |
| } | |
| }, { | |
| full_name: "hal.command_buffer.execution_barrier", | |
| signature: { | |
| argument_types: [ 3, 0, 0, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.command_buffer.fill_buffer", | |
| signature: { | |
| argument_types: [ 3, 2, 0, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.command_buffer.copy_buffer", | |
| signature: { | |
| argument_types: [ 3, 2, 0, 2, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.command_buffer.bind_descriptor_set", | |
| signature: { | |
| argument_types: [ 3, 5, 0, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.command_buffer.dispatch", | |
| signature: { | |
| argument_types: [ 3, 5, 0, 0, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.command_buffer.dispatch.indirect", | |
| signature: { | |
| argument_types: [ 3, 5, 0, 2, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.descriptor_set.allocate", | |
| signature: { | |
| argument_types: [ 4, 0 ], | |
| result_types: [ 0 ] | |
| } | |
| }, { | |
| full_name: "hal.descriptor_set.update", | |
| signature: { | |
| argument_types: [ 4, 0, 0, 2, 0, 0, 0 ] | |
| } | |
| }, { | |
| full_name: "hal.device.allocator", | |
| signature: { | |
| argument_types: [ 4 ], | |
| result_types: [ 1 ] | |
| } | |
| } ], | |
| exported_functions: [ { | |
| local_name: "hloElementwiseOps", | |
| signature: { | |
| argument_types: [ 2 ], | |
| result_types: [ 2 ] | |
| }, | |
| internal_ordinal: 1 | |
| } ], | |
| internal_functions: [ { | |
| local_name: "hloElementwiseOps_ex_dispatch_0", | |
| signature: { | |
| argument_types: [ 4 ], | |
| result_types: [ 5 ] | |
| } | |
| }, { | |
| local_name: "hloElementwiseOps", | |
| signature: { | |
| argument_types: [ 2 ], | |
| result_types: [ 2 ], | |
| reflection_attrs: [ { | |
| key: "f", | |
| value: "I6!B3!d4R6!B3!d4" | |
| }, { | |
| key: "fv", | |
| value: "1" | |
| } ] | |
| } | |
| } ], | |
| rodata_segments: [ { | |
| data: [ 16, 0, 0, 0, 69, 77, 79, 68, 8, 0, 14, 0, 8, 0, 4, 0, 8, 0, 0, 0, 20, 0, 0, 0, 52, 0, 0, 0, 0, 0, 10, 0, 16, 0, 12, 0, 8, 0, 4, 0, 10, 0, 0, 0, 12, 0, 0, 0, 16, 0, 0, 0, 16, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 28, 0, 0, 0, 6, 0, 0, 0, 109, 111, 100, 117, 108, 101, 0, 0, 12, 0, 16, 0, 12, 0, 8, 0, 0, 0, 4, 0, 12, 0, 0, 0, 204, 0, 0, 0, 48, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 126, 255, 255, 255, 4, 0, 0, 0, 2, 0, 0, 0, 68, 0, 0, 0, 4, 0, 0, 0, 204, 255, 255, 255, 8, 0, 0, 0, 0, 0, 0, 1, 144, 255, 255, 255, 8, 0, 0, 0, 12, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, 236, 255, 255, 255, 8, 0, 0, 0, 0, 0, 0, 1, 190, 255, 255, 255, 32, 0, 0, 0, 8, 0, 12, 0, 11, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 1, 204, 255, 255, 255, 8, 0, 0, 0, 20, 0, 0, 0, 1, 0, 0, 0, 4, 0, 0, 0, 8, 0, 10, 0, 9, 0, 4, 0, 8, 0, 0, 0, 12, 0, 0, 0, 0, 1, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 32, 0, 0, 0, 8, 0, 12, 0, 8, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 7, 0, 0, 0, 112, 0, 0, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 2, 0, 113, 0, 0, 0, 0, 2, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 3, 0, 115, 2, 0, 0, 0, 3, 0, 35, 0, 0, 0, 0, 5, 1, 4, 0, 0, 0, 0, 4, 0, 119, 3, 0, 0, 0, 4, 0, 0, 3, 1, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 5, 0, 0, 3, 1, 1, 0, 0, 0, 1, 4, 0, 0, 0, 0, 0, 0, 0, 6, 0, 54, 4, 0, 5, 0, 1, 0, 5, 0, 6, 0, 4, 0 ] | |
| }, { | |
| data: [ 8, 0, 0, 0, 83, 80, 86, 69, 164, 255, 255, 255, 104, 4, 0, 0, 112, 0, 0, 0, 4, 0, 0, 0, 238, 255, 255, 255, 4, 0, 0, 0, 1, 0, 0, 0, 12, 0, 0, 0, 0, 0, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 52, 0, 0, 0, 16, 0, 0, 0, 12, 0, 20, 0, 4, 0, 8, 0, 12, 0, 16, 0, 12, 0, 0, 0, 1, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, 12, 0, 16, 0, 0, 0, 4, 0, 8, 0, 12, 0, 12, 0, 0, 0, 7, 0, 0, 0, 1, 0, 0, 0, 32, 0, 0, 0, 252, 0, 0, 0, 3, 2, 35, 7, 0, 0, 1, 0, 22, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 17, 0, 2, 0, 1, 0, 0, 0, 10, 0, 11, 0, 83, 80, 86, 95, 75, 72, 82, 95, 115, 116, 111, 114, 97, 103, 101, 95, 98, 117, 102, 102, 101, 114, 95, 115, 116, 111, 114, 97, 103, 101, 95, 99, 108, 97, 115, 115, 0, 0, 0, 0, 14, 0, 3, 0, 0, 0, 0, 0, 1, 0, 0, 0, 15, 0, 13, 0, 5, 0, 0, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 4, 0, 0, 0, 16, 0, 6, 0, 14, 0, 0, 0, 17, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 5, 0, 7, 0, 4, 0, 0, 0, 103, 108, 111, 98, 97, 108, 73, 110, 118, 111, 99, 97, 116, 105, 111, 110, 73, 68, 0, 0, 5, 0, 12, 0, 10, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 48, 0, 0, 5, 0, 12, 0, 11, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 95, 97, 114, 103, 95, 49, 0, 0, 5, 0, 11, 0, 14, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 71, 0, 4, 0, 4, 0, 0, 0, 11, 0, 0, 0, 28, 0, 0, 0, 71, 0, 4, 0, 7, 0, 0, 0, 6, 0, 0, 0, 4, 0, 0, 0, 72, 0, 5, 0, 6, 0, 0, 0, 0, 0, 0, 0, 35, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 33, 0, 0, 0, 0, 0, 0, 0, 71, 0, 4, 0, 10, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 71, 0, 3, 0, 6, 0, 0, 0, 2, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 33, 0, 0, 0, 1, 0, 0, 0, 71, 0, 4, 0, 11, 0, 0, 0, 34, 0, 0, 0, 0, 0, 0, 0, 21, 0, 4, 0, 3, 0, 0, 0, 32, 0, 0, 0, 1, 0, 0, 0, 23, 0, 4, 0, 2, 0, 0, 0, 3, 0, 0, 0, 3, 0, 0, 0, 32, 0, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 2, 0, 0, 0, 59, 0, 4, 0, 1, 0, 0, 0, 4, 0, 0, 0, 1, 0, 0, 0, 22, 0, 3, 0, 8, 0, 0, 0, 32, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 9, 0, 0, 0, 4, 0, 0, 0, 28, 0, 4, 0, 7, 0, 0, 0, 8, 0, 0, 0, 9, 0, 0, 0, 30, 0, 3, 0, 6, 0, 0, 0, 7, 0, 0, 0, 32, 0, 4, 0, 5, 0, 0, 0, 12, 0, 0, 0, 6, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 10, 0, 0, 0, 12, 0, 0, 0, 59, 0, 4, 0, 5, 0, 0, 0, 11, 0, 0, 0, 12, 0, 0, 0, 19, 0, 2, 0, 13, 0, 0, 0, 33, 0, 3, 0, 12, 0, 0, 0, 13, 0, 0, 0, 20, 0, 2, 0, 16, 0, 0, 0, 41, 0, 3, 0, 16, 0, 0, 0, 17, 0, 0, 0, 43, 0, 4, 0, 3, 0, 0, 0, 18, 0, 0, 0, 0, 0, 0, 0, 32, 0, 4, 0, 26, 0, 0, 0, 12, 0, 0, 0, 8, 0, 0, 0, 54, 0, 5, 0, 13, 0, 0, 0, 14, 0, 0, 0, 0, 0, 0, 0, 12, 0, 0, 0, 248, 0, 2, 0, 15, 0, 0, 0, 61, 0, 4, 0, 2, 0, 0, 0, 19, 0, 0, 0, 4, 0, 0, 0, 81, 0, 5, 0, 3, 0, 0, 0, 20, 0, 0, 0, 19, 0, 0, 0, 0, 0, 0, 0, 177, 0, 5, 0, 16, 0, 0, 0, 21, 0, 0, 0, 20, 0, 0, 0, 9, 0, 0, 0, 167, 0, 5, 0, 16, 0, 0, 0, 22, 0, 0, 0, 21, 0, 0, 0, 17, 0, 0, 0, 247, 0, 3, 0, 25, 0, 0, 0, 0, 0, 0, 0, 250, 0, 4, 0, 22, 0, 0, 0, 24, 0, 0, 0, 25, 0, 0, 0, 248, 0, 2, 0, 24, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 27, 0, 0, 0, 10, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 61, 0, 4, 0, 8, 0, 0, 0, 28, 0, 0, 0, 27, 0, 0, 0, 129, 0, 5, 0, 8, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, 28, 0, 0, 0, 131, 0, 5, 0, 8, 0, 0, 0, 30, 0, 0, 0, 29, 0, 0, 0, 28, 0, 0, 0, 133, 0, 5, 0, 8, 0, 0, 0, 31, 0, 0, 0, 30, 0, 0, 0, 28, 0, 0, 0, 65, 0, 6, 0, 26, 0, 0, 0, 32, 0, 0, 0, 11, 0, 0, 0, 18, 0, 0, 0, 20, 0, 0, 0, 62, 0, 3, 0, 32, 0, 0, 0, 31, 0, 0, 0, 249, 0, 2, 0, 25, 0, 0, 0, 248, 0, 2, 0, 25, 0, 0, 0, 253, 0, 1, 0, 56, 0, 1, 0, 1, 0, 0, 0, 4, 0, 0, 0, 32, 0, 0, 0, 104, 108, 111, 69, 108, 101, 109, 101, 110, 116, 119, 105, 115, 101, 79, 112, 115, 95, 114, 103, 110, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0 ] | |
| } ], | |
| module_state: { | |
| global_ref_count: 1 | |
| }, | |
| function_descriptors: [ { | |
| bytecode_offset: 0, | |
| bytecode_length: 162, | |
| i32_register_count: 4, | |
| ref_register_count: 4 | |
| }, { | |
| bytecode_offset: 162, | |
| bytecode_length: 377, | |
| i32_register_count: 11, | |
| ref_register_count: 7 | |
| } ], | |
| bytecode_data: [ 9, 69, 69, 82, 73, 0, 9, 69, 86, 80, 83, 1, 10, 129, 2, 0, 0, 0, 0, 5, 0, 0, 0, 130, 76, 130, 2, 81, 2, 41, 0, 0, 0, 1, 130, 128, 44, 0, 0, 0, 0, 84, 1, 192, 83, 1, 0, 0, 128, 2, 255, 2, 3, 128, 0, 1, 1, 2, 80, 64, 0, 0, 0, 0, 64, 2, 0, 0, 81, 0, 84, 0, 0, 0, 1, 2, 0, 107, 0, 0, 0, 1, 2, 0, 11, 0, 0, 0, 0, 129, 82, 2, 0, 0, 128, 3, 192, 0, 193, 1, 128, 80, 146, 0, 0, 0, 0, 64, 0, 1, 1, 81, 1, 123, 0, 0, 0, 0, 159, 0, 0, 0, 0, 11, 1, 0, 0, 0, 129, 82, 2, 0, 0, 128, 3, 192, 0, 193, 1, 128, 80, 146, 0, 0, 0, 0, 3, 0, 0, 0, 0, 5, 0, 0, 0, 128, 84, 1, 192, 84, 1, 193, 9, 1, 0, 0, 0, 0, 9, 4, 0, 0, 0, 1, 9, 50, 0, 0, 0, 2, 9, 15, 0, 0, 0, 3, 9, 3, 0, 0, 0, 4, 9, 69, 69, 82, 73, 5, 9, 69, 86, 80, 83, 6, 10, 129, 8, 7, 9, 16, 0, 0, 0, 8, 9, 8, 0, 0, 0, 9, 82, 0, 0, 0, 128, 0, 1, 130, 82, 33, 0, 0, 128, 1, 130, 1, 131, 83, 10, 0, 0, 128, 5, 255, 255, 255, 1, 255, 5, 195, 2, 3, 1, 1, 1, 131, 82, 5, 0, 0, 128, 1, 131, 0, 82, 22, 0, 0, 128, 3, 130, 0, 4, 1, 132, 82, 23, 0, 0, 128, 1, 132, 0, 2, 0, 0, 0, 0, 5, 0, 0, 0, 133, 76, 133, 2, 81, 2, 148, 0, 0, 0, 1, 133, 129, 154, 0, 0, 0, 0, 80, 18, 1, 0, 0, 0, 83, 1, 0, 0, 128, 2, 255, 2, 3, 130, 5, 6, 1, 2, 80, 174, 0, 0, 0, 0, 64, 2, 5, 3, 81, 3, 190, 0, 0, 0, 0, 213, 0, 0, 0, 0, 11, 0, 0, 0, 0, 129, 82, 2, 0, 0, 128, 3, 130, 2, 193, 1, 129, 80, 252, 0, 0, 0, 0, 64, 2, 6, 3, 81, 3, 229, 0, 0, 0, 0, 12, 1, 0, 0, 0, 11, 1, 0, 0, 0, 129, 82, 2, 0, 0, 128, 3, 130, 2, 193, 1, 129, 80, 252, 0, 0, 0, 0, 3, 0, 0, 0, 0, 5, 0, 0, 0, 129, 80, 18, 1, 0, 0, 0, 80, 18, 1, 0, 0, 0, 83, 3, 0, 0, 128, 5, 255, 255, 255, 1, 255, 5, 132, 7, 128, 1, 1, 0, 82, 5, 0, 0, 128, 1, 192, 0, 83, 3, 0, 0, 128, 5, 255, 255, 255, 1, 255, 5, 132, 0, 131, 1, 1, 0, 82, 5, 0, 0, 128, 1, 131, 0, 82, 29, 0, 0, 128, 6, 132, 193, 7, 0, 0, 0, 0, 83, 25, 0, 0, 128, 5, 255, 255, 255, 1, 0, 5, 132, 8, 0, 9, 1, 0, 82, 24, 0, 0, 128, 1, 132, 0, 82, 6, 0, 0, 128, 2, 194, 196, 0, 84, 1, 195 ] | |
| } |
This file contains 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
| bazel build iree/iree/tools:iree-translate | |
| ./bazel-bin/third_party/iree/iree/tools/iree-translate -print-ir-after-all -split-input-file -iree-mlir-to-vm-bytecode-module -iree-vm-bytecode-module-output-format=flatbuffer-text iree/iree/compiler/Translation/test/smoketest.mlir |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment