iree-org / iree

A retargetable MLIR-based machine learning compiler and runtime toolkit.
http://iree.dev/
Apache License 2.0
2.79k stars 603 forks source link

iree-import-xla fails to lower lax.while_loop #5648

Closed phoenix-meadowlark closed 2 years ago

phoenix-meadowlark commented 3 years ago

Lowering lax.while_loop fails due to error: 'mhlo.constant' op operation destroyed but still has uses. The python, HLO text and stack dump / reproduction instructions can be found below.

def cond_fun(state):
  i, _ = state
  return i < 8

def body_fun(state):
  i, array = state
  return (i + 1, array + i)

def while_loop(x):
  return lax.while_loop(cond_fun, body_fun, (0, x))
HloModule xla_computation_wrapped_function__10.30

body_computation__3.5 {
  constant.9 = pred[] constant(false)
  parameter.6 = (s32[], s32[3]{0}) parameter(0)
  get-tuple-element.7 = s32[] get-tuple-element(parameter.6), index=0
  constant.10 = s32[] constant(1)
  add.11 = s32[] add(get-tuple-element.7, constant.10)
  get-tuple-element.8 = s32[3]{0} get-tuple-element(parameter.6), index=1
  convert.12 = s32[] convert(get-tuple-element.7)
  broadcast.13 = s32[3]{0} broadcast(convert.12), dimensions={}
  add.14 = s32[3]{0} add(get-tuple-element.8, broadcast.13)
  ROOT tuple.15 = (s32[], s32[3]{0}) tuple(add.11, add.14)
}

cond_computation__3.16 {
  parameter.17 = (s32[], s32[3]{0}) parameter(0)
  get-tuple-element.19 = s32[3]{0} get-tuple-element(parameter.17), index=1
  constant.20 = pred[] constant(false)
  get-tuple-element.18 = s32[] get-tuple-element(parameter.17), index=0
  constant.21 = s32[] constant(8)
  ROOT compare.22 = pred[] compare(get-tuple-element.18, constant.21), direction=LT
}

ENTRY xla_computation_wrapped_function__10.30 {
  constant.2 = pred[] constant(false)
  constant.3 = s32[] constant(0)
  parameter.1 = s32[3]{0} parameter(0)
  tuple.4 = (s32[], s32[3]{0}) tuple(constant.3, parameter.1)
  while.23 = (s32[], s32[3]{0}) while(tuple.4), condition=cond_computation__3.16, body=body_computation__3.5
  get-tuple-element.24 = s32[] get-tuple-element(while.23), index=0
  get-tuple-element.25 = s32[3]{0} get-tuple-element(while.23), index=1
  tuple.26 = (s32[], s32[3]{0}) tuple(get-tuple-element.24, get-tuple-element.25)
  get-tuple-element.27 = s32[] get-tuple-element(tuple.26), index=0
  get-tuple-element.28 = s32[3]{0} get-tuple-element(tuple.26), index=1
  ROOT tuple.29 = (s32[], s32[3]{0}) tuple(get-tuple-element.27, get-tuple-element.28)
}
iree-import-xla --xla-format=hlo_text /tmp/lax_while.hlo.txt
LLVM ERROR: operation destroyed but still has uses
Stack dump:
0.  Program arguments: iree-import-xla --xla-format=hlo_text /tmp/lax_while.hlo.txt
1.  In-Flight Diagnostics:
    loc("constant.21"): error: 'mhlo.constant' op operation destroyed but still has uses
 #0 0x0000557b2fa5d784 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) /proc/self/cwd/third_party/llvm/llvm-project/llvm/lib/Support/Unix/Signals.inc:570:11
 #1 0x0000557b2fa5dbeb PrintStackTraceSignalHandler(void*) /proc/self/cwd/third_party/llvm/llvm-project/llvm/lib/Support/Unix/Signals.inc:632:1
 #2 0x0000557b2fa5c0e3 llvm::sys::RunSignalHandlers() /proc/self/cwd/third_party/llvm/llvm-project/llvm/lib/Support/Signals.cpp:76:5
 #3 0x0000557b2fa5e215 SignalHandler(int) /proc/self/cwd/third_party/llvm/llvm-project/llvm/lib/Support/Unix/Signals.inc:407:1
 #4 0x00007fbb1ffc99a0 __restore_rt (/usr/grte/v4/lib64/libpthread.so.0+0xf9a0)
 #5 0x00007fbb1fe4b6c2 raise (/usr/grte/v4/lib64/libc.so.6+0x4c6c2)
 #6 0x00007fbb1fe4d3e0 abort (/usr/grte/v4/lib64/libc.so.6+0x4e3e0)
 #7 0x0000557b2fa2c032 llvm::report_fatal_error(llvm::Twine const&, bool) /proc/self/cwd/third_party/llvm/llvm-project/llvm/lib/Support/ErrorHandling.cpp:126:3
 #8 0x0000557b2fa2bea2 llvm::report_fatal_error(char const*, bool) /proc/self/cwd/third_party/llvm/llvm-project/llvm/lib/Support/ErrorHandling.cpp:83:3
 #9 0x0000557b2f99bf7e mlir::Operation::~Operation() /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp:193:5
#10 0x0000557b2f99c21d mlir::Operation::destroy() /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp:216:8
#11 0x0000557b2f99ce45 llvm::ilist_traits<mlir::Operation>::deleteNode(mlir::Operation*) /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp:440:1
#12 0x0000557b2f8c6145 llvm::iplist_impl<llvm::simple_ilist<mlir::Operation>, llvm::ilist_traits<mlir::Operation> >::erase(llvm::ilist_iterator<llvm::ilist_detail::node_options<mlir::Operation, false, false, void>, false, false>) /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/ilist.h:269:12
#13 0x0000557b2f8c610e llvm::iplist_impl<llvm::simple_ilist<mlir::Operation>, llvm::ilist_traits<mlir::Operation> >::pop_back() /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/ilist.h:320:25
#14 0x0000557b2f8c53be mlir::Block::clear() /proc/self/cwd/third_party/llvm/llvm-project/mlir/include/mlir/IR/Block.h:40:5
#15 0x0000557b2f8c3abd mlir::Block::~Block() /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/IR/Block.cpp:21:3
#16 0x0000557b2a960b07 llvm::ilist_alloc_traits<mlir::Block>::deleteNode(mlir::Block*) /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/ilist.h:41:39
#17 0x0000557b2a960ac5 llvm::iplist_impl<llvm::simple_ilist<mlir::Block>, llvm::ilist_traits<mlir::Block> >::erase(llvm::ilist_iterator<llvm::ilist_detail::node_options<mlir::Block, false, false, void>, false, false>) /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/ilist.h:269:12
#18 0x0000557b2a960a7b llvm::iplist_impl<llvm::simple_ilist<mlir::Block>, llvm::ilist_traits<mlir::Block> >::erase(llvm::ilist_iterator<llvm::ilist_detail::node_options<mlir::Block, false, false, void>, false, false>, llvm::ilist_iterator<llvm::ilist_detail::node_options<mlir::Block, false, false, void>, false, false>) /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/ilist.h:305:15
#19 0x0000557b2a96023b llvm::iplist_impl<llvm::simple_ilist<mlir::Block>, llvm::ilist_traits<mlir::Block> >::clear() /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/ilist.h:309:18
#20 0x0000557b2f9b2275 llvm::iplist_impl<llvm::simple_ilist<mlir::Block>, llvm::ilist_traits<mlir::Block> >::~iplist_impl() /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/ilist.h:210:29
#21 0x0000557b2f9b20c5 llvm::iplist<mlir::Block>::~iplist() /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/ilist.h:390:7
#22 0x0000557b2f9b0a52 mlir::Region::~Region() /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/IR/Region.cpp:20:1
#23 0x0000557b2f99c0c8 mlir::Operation::~Operation() /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp:205:21
#24 0x0000557b2f99c21d mlir::Operation::destroy() /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp:216:8
#25 0x0000557b2f99ce45 llvm::ilist_traits<mlir::Operation>::deleteNode(mlir::Operation*) /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp:440:1
#26 0x0000557b2f8c6145 llvm::iplist_impl<llvm::simple_ilist<mlir::Operation>, llvm::ilist_traits<mlir::Operation> >::erase(llvm::ilist_iterator<llvm::ilist_detail::node_options<mlir::Operation, false, false, void>, false, false>) /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/ilist.h:269:12
#27 0x0000557b2f9a42a2 llvm::iplist_impl<llvm::simple_ilist<mlir::Operation>, llvm::ilist_traits<mlir::Operation> >::erase(mlir::Operation*) /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/ilist.h:272:39
#28 0x0000557b2f99d07d mlir::Operation::erase() /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp:488:29
#29 0x0000557b2a94bf48 mlir::OpState::erase() /proc/self/cwd/third_party/llvm/llvm-project/mlir/include/mlir/IR/OpDefinition.h:130:34
#30 0x0000557b2b036969 mlir::mhlo::(anonymous namespace)::MatchAndRewrite(mlir::mhlo::WhileOp) /proc/self/cwd/third_party/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/mhlo_control_flow_to_scf.cc:192:1
#31 0x0000557b2b035891 mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction()::'lambda'(mlir::mhlo::WhileOp)::operator()(mlir::mhlo::WhileOp) const /proc/self/cwd/third_party/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/mhlo_control_flow_to_scf.cc:47:73
#32 0x0000557b2b035863 std::__u::enable_if<(!(llvm::is_one_of<mlir::mhlo::WhileOp, mlir::Operation*, mlir::Region*, mlir::Block*>::value)) && (std::is_same<void, void>::value), void>::type mlir::detail::walk<(mlir::WalkOrder)1, mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction()::'lambda'(mlir::mhlo::WhileOp), mlir::mhlo::WhileOp, void>(mlir::Operation*, mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction()::'lambda'(mlir::mhlo::WhileOp)&&)::'lambda'(mlir::Operation*)::operator()(mlir::Operation*) const /proc/self/cwd/third_party/llvm/llvm-project/mlir/include/mlir/IR/Visitors.h:160:3
#33 0x0000557b2b0357fd void llvm::function_ref<void (mlir::Operation*)>::callback_fn<std::__u::enable_if<(!(llvm::is_one_of<mlir::mhlo::WhileOp, mlir::Operation*, mlir::Region*, mlir::Block*>::value)) && (std::is_same<void, void>::value), void>::type mlir::detail::walk<(mlir::WalkOrder)1, mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction()::'lambda'(mlir::mhlo::WhileOp), mlir::mhlo::WhileOp, void>(mlir::Operation*, mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction()::'lambda'(mlir::mhlo::WhileOp)&&)::'lambda'(mlir::Operation*)>(long, mlir::Operation*) /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLExtras.h:185:5
#34 0x0000557b2f0be34c llvm::function_ref<void (mlir::Operation*)>::operator()(mlir::Operation*) const /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLExtras.h:209:5
#35 0x0000557b2f9cfe69 mlir::detail::walk(mlir::Operation*, llvm::function_ref<void (mlir::Operation*)>, mlir::WalkOrder) /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/IR/Visitors.cpp:68:1
#36 0x0000557b2f9cfe18 mlir::detail::walk(mlir::Operation*, llvm::function_ref<void (mlir::Operation*)>, mlir::WalkOrder) /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/IR/Visitors.cpp:61:27
#37 0x0000557b2b03579a std::__u::enable_if<(!(llvm::is_one_of<mlir::mhlo::WhileOp, mlir::Operation*, mlir::Region*, mlir::Block*>::value)) && (std::is_same<void, void>::value), void>::type mlir::detail::walk<(mlir::WalkOrder)1, mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction()::'lambda'(mlir::mhlo::WhileOp), mlir::mhlo::WhileOp, void>(mlir::Operation*, mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction()::'lambda'(mlir::mhlo::WhileOp)&&) /proc/self/cwd/third_party/llvm/llvm-project/mlir/include/mlir/IR/Visitors.h:161:3
#38 0x0000557b2b03572d void mlir::Operation::walk<(mlir::WalkOrder)1, mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction()::'lambda'(mlir::mhlo::WhileOp), void>(mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction()::'lambda'(mlir::mhlo::WhileOp)&&) /proc/self/cwd/third_party/llvm/llvm-project/mlir/include/mlir/IR/Operation.h:522:5
#39 0x0000557b2b0356f0 void mlir::OpState::walk<(mlir::WalkOrder)1, mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction()::'lambda'(mlir::mhlo::WhileOp), void>(mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction()::'lambda'(mlir::mhlo::WhileOp)&&) /proc/self/cwd/third_party/llvm/llvm-project/mlir/include/mlir/IR/OpDefinition.h:161:5
#40 0x0000557b2b0351b6 mlir::mhlo::(anonymous namespace)::ControlFlowToScfPass::runOnFunction() /proc/self/cwd/third_party/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/mhlo_control_flow_to_scf.cc:48:3
#41 0x0000557b2a93b4cf mlir::FunctionPass::runOnOperation() /proc/self/cwd/third_party/llvm/llvm-project/mlir/include/mlir/Pass/Pass.h:379:3
#42 0x0000557b2f78ecd2 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:407:21
#43 0x0000557b2f78f255 mlir::detail::OpToOpPassAdaptor::runPipeline(llvm::iterator_range<llvm::pointee_iterator<std::__u::unique_ptr<mlir::Pass, std::__u::default_delete<mlir::Pass> >*, mlir::Pass> >, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:448:16
#44 0x0000557b2f795656 mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_8::operator()(llvm::MutableArrayRef<mlir::OpPassManager>) const /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:623:15
#45 0x0000557b2f794fac mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_8 std::__u::for_each<llvm::SmallVector<mlir::OpPassManager, 1u>*, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_8>(llvm::SmallVector<mlir::OpPassManager, 1u>*, llvm::SmallVector<mlir::OpPassManager, 1u>*, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_8) /proc/self/cwd/third_party/crosstool/v18/stable/toolchain/bin/../include/c++/v1/algorithm:895:31
#46 0x0000557b2f794f47 void llvm::parallel::detail::parallel_for_each<llvm::SmallVector<mlir::OpPassManager, 1u>*, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_8>(llvm::SmallVector<mlir::OpPassManager, 1u>*, llvm::SmallVector<mlir::OpPassManager, 1u>*, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_8) /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/Support/Parallel.h:144:1
#47 0x0000557b2f79034e void llvm::parallelForEach<llvm::SmallVector<mlir::OpPassManager, 1u>*, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_8>(llvm::SmallVector<mlir::OpPassManager, 1u>*, llvm::SmallVector<mlir::OpPassManager, 1u>*, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_8) /proc/self/cwd/third_party/llvm/llvm-project/llvm/include/llvm/Support/Parallel.h:234:5
#48 0x0000557b2f78fca8 mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool) /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:641:7
#49 0x0000557b2f78ef27 mlir::detail::OpToOpPassAdaptor::runOnOperation(bool) /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:523:5
#50 0x0000557b2f78ecc3 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:404:5
#51 0x0000557b2f78f255 mlir::detail::OpToOpPassAdaptor::runPipeline(llvm::iterator_range<llvm::pointee_iterator<std::__u::unique_ptr<mlir::Pass, std::__u::default_delete<mlir::Pass> >*, mlir::Pass> >, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:448:16
#52 0x0000557b2f791104 mlir::PassManager::run(mlir::Operation*) /proc/self/cwd/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:900:13
#53 0x0000557b2a91fcf4 main /proc/self/cwd/third_party/iree/integrations/tensorflow/iree_tf_compiler/iree-import-xla-main.cpp:229:17
#54 0x00007fbb1fe37bbd __libc_start_main (/usr/grte/v4/lib64/libc.so.6+0x38bbd)
#55 0x0000557b2a91f029 _start /usr/grte/v4/debug-src/src/csu/../sysdeps/x86_64/start.S:121:0
[1]    662240 abort       --xla-format=hlo_text /tmp/lax_while.hlo.txt
phoenix-meadowlark commented 3 years ago

It appears that lax.fori_loop gives the same error when given a parameter (in this case an array of three integers):

def add_45(x):

  def add(i, x):
    return x + i

  x = lax.fori_loop(0, 10, add, x)
  return x
HloModule xla_computation_wrapped_function__1.41

body_computation__1.7 {
  constant.12 = pred[] constant(false)
  parameter.8 = (s32[], s32[], s32[3]{0}) parameter(0)
  get-tuple-element.9 = s32[] get-tuple-element(parameter.8), index=0
  constant.18 = s32[] constant(1)
  add.19 = s32[] add(get-tuple-element.9, constant.18)
  get-tuple-element.10 = s32[] get-tuple-element(parameter.8), index=1
  constant.13 = s32[] constant(1)
  add.14 = s32[] add(get-tuple-element.10, constant.13)
  convert.15 = s32[] convert(get-tuple-element.10)
  broadcast.16 = s32[3]{0} broadcast(convert.15), dimensions={}
  get-tuple-element.11 = s32[3]{0} get-tuple-element(parameter.8), index=2
  add.17 = s32[3]{0} add(broadcast.16, get-tuple-element.11)
  ROOT tuple.20 = (s32[], s32[], s32[3]{0}) tuple(add.19, add.14, add.17)
}

cond_computation__1.21 {
  parameter.22 = (s32[], s32[], s32[3]{0}) parameter(0)
  get-tuple-element.24 = s32[] get-tuple-element(parameter.22), index=1
  get-tuple-element.25 = s32[3]{0} get-tuple-element(parameter.22), index=2
  constant.26 = pred[] constant(false)
  get-tuple-element.23 = s32[] get-tuple-element(parameter.22), index=0
  constant.27 = s32[] constant(10)
  ROOT compare.28 = pred[] compare(get-tuple-element.23, constant.27), direction=LT
}

ENTRY xla_computation_wrapped_function__1.41 {
  constant.2 = pred[] constant(false)
  constant.4 = pred[] constant(false)
  constant.5 = s32[] constant(0)
  constant.3 = s32[] constant(0)
  parameter.1 = s32[3]{0} parameter(0)
  tuple.6 = (s32[], s32[], s32[3]{0}) tuple(constant.5, constant.3, parameter.1)
  while.29 = (s32[], s32[], s32[3]{0}) while(tuple.6), condition=cond_computation__1.21, body=body_computation__1.7
  get-tuple-element.30 = s32[] get-tuple-element(while.29), index=0
  get-tuple-element.31 = s32[] get-tuple-element(while.29), index=1
  get-tuple-element.32 = s32[3]{0} get-tuple-element(while.29), index=2
  tuple.33 = (s32[], s32[], s32[3]{0}) tuple(get-tuple-element.30, get-tuple-element.31, get-tuple-element.32)
  get-tuple-element.34 = s32[] get-tuple-element(tuple.33), index=0
  get-tuple-element.35 = s32[] get-tuple-element(tuple.33), index=1
  get-tuple-element.36 = s32[3]{0} get-tuple-element(tuple.33), index=2
  tuple.37 = (s32[], s32[3]{0}) tuple(get-tuple-element.35, get-tuple-element.36)
  get-tuple-element.38 = s32[] get-tuple-element(tuple.37), index=0
  get-tuple-element.39 = s32[3]{0} get-tuple-element(tuple.37), index=1
  ROOT tuple.40 = (s32[3]{0}) tuple(get-tuple-element.39)
}
LLVM ERROR: operation destroyed but still has uses
Stack dump:
0.  Program arguments: iree-import-xla --xla-format=hlo_text /tmp/lax_fori_loop.hlo.txt
1.  In-Flight Diagnostics:
    loc("constant.27"): error: 'mhlo.constant' op operation destroyed but still has uses
 ...
[1]    695151 abort       --xla-format=hlo_text /tmp/lax_fori_loop.hlo.txt

That seems reasonable since they just call the same op. The function I posted in #5510 still produces the operand #0 does not dominate this use error though, even though it should be identical as far as the loop is concerned. (It looks like the difference in number of loop iterations also isn't the cause here).

stellaraccident commented 3 years ago

With the latest updated, this now compiles but fails at runtime with:

ValueError: Error invoking function: third_party/iree/iree/modules/hal/hal_module.c:362: INVALID_ARGUMENT; load length byte count 16 exceeds max; while invoking native function hal.buffer.load; while calling import
phoenix-meadowlark commented 2 years ago

Closing this in favor of #9499 as the HLO importer is defunct.