ireec crash for mhlo.pad operator

[hawkinsp]

import iree.compiler

CODE = """
#loc0 = loc(unknown)
module @jit_prim_fun.0 {
  func.func public @main(%arg0: tensor<2x3xi8> loc(unknown), %arg1: tensor<i8> loc(unknown)) -> tensor<2x1xi8> {
    %0 = "mhlo.pad"(%arg0, %arg1) {edge_padding_high = dense<[0, -1]> : tensor<2xi64>, edge_padding_low = dense<[0, -1]> : tensor<2xi64>, interior_padding = dense<0> : tensor<2xi64>} : (tensor<2x3xi8>, tensor<i8>) -> tensor<2x1xi8> loc(#loc1)
    return %0 : tensor<2x1xi8> loc(#loc0)
  } loc(#loc0)
} loc(#loc0)
#loc1 = loc("jit(pad)/jit(main)/pad[padding_config=((0, 0, 0), (-1, -1, 0))]"("/Users/phawkins/p/jax/tests/lax_test.py":1439:1))
"""

iree_binary = iree.compiler.compile_str(
    CODE, target_backends=["dylib"], input_type="mhlo")

Output:

Traceback (most recent call last):
  File "/Users/phawkins/p/jax/t.py", line 14, in <module>
    iree_binary = iree.compiler.compile_str(
  File "/Users/phawkins/.pyenv/versions/py310/lib/python3.10/site-packages/iree/compiler/tools/core.py", line 293, in compile_str
    result = invoke_immediate(cl, immediate_input=input_bytes)
  File "/Users/phawkins/.pyenv/versions/py310/lib/python3.10/site-packages/iree/compiler/tools/binaries.py", line 200, in invoke_immediate
    raise CompilerToolError(process)
iree.compiler.tools.binaries.CompilerToolError: Error invoking IREE compiler tool iree-compile
Diagnostics:
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.	Program arguments: /Users/phawkins/.pyenv/versions/py310/lib/python3.10/site-packages/iree/compiler/tools/../_mlir_libs/iree-compile - --iree-input-type=mhlo --iree-vm-bytecode-module-output-format=flatbuffer-binary --iree-hal-target-backends=dylib --iree-mlir-to-vm-bytecode-module --iree-llvm-embedded-linker-path=/Users/phawkins/.pyenv/versions/py310/lib/python3.10/site-packages/iree/compiler/tools/../_mlir_libs/iree-lld --mlir-print-debuginfo --mlir-print-op-on-diagnostic=false
Stack dump without symbol names (ensure you have llvm-symbolizer in your PATH or set the environment var `LLVM_SYMBOLIZER_PATH` to point to it):
0  libIREECompilerAggregateCAPI.dylib 0x000000010da184ac llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) + 56
1  libIREECompilerAggregateCAPI.dylib 0x000000010da1748c llvm::sys::RunSignalHandlers() + 128
2  libIREECompilerAggregateCAPI.dylib 0x000000010da18b00 llvm::sys::PrintStackTraceOnErrorSignal(llvm::StringRef, bool) + 708
3  libsystem_platform.dylib           0x00000001b06534a4 _sigtramp + 56
4  libsystem_platform.dylib           0xffff80010daadb34 _sigtramp + 18446603333491074760
5  libIREECompilerAggregateCAPI.dylib 0x000000010efafd94 spirv_cross::DominatorBuilder::lift_continue_block_dominator() + 60692
6  libIREECompilerAggregateCAPI.dylib 0x00000001110fcdf4 mlir::getConstantIntValue(mlir::OpFoldResult) + 80
7  libIREECompilerAggregateCAPI.dylib 0x000000010e548f0c mlir::tensor::YieldOp::getEffects(llvm::SmallVectorImpl<mlir::SideEffects::EffectInstance<mlir::MemoryEffects::Effect>>&) + 38076
8  libIREECompilerAggregateCAPI.dylib 0x000000010de8e9c4 mlir::createLinalgTilingPass(llvm::ArrayRef<long long>, mlir::linalg::LinalgTilingLoopType) + 7632
9  libIREECompilerAggregateCAPI.dylib 0x0000000111165640 mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<mlir::LogicalResult (mlir::Pattern const&)>) + 1196
10 libIREECompilerAggregateCAPI.dylib 0x0000000111146d8c mlir::applyPatternsAndFoldGreedily(llvm::MutableArrayRef<mlir::Region>, mlir::FrozenRewritePatternSet const&, mlir::GreedyRewriteConfig) + 1520
11 libIREECompilerAggregateCAPI.dylib 0x000000010da4a384 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) + 724
12 libIREECompilerAggregateCAPI.dylib 0x000000010da4ac3c mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) + 348
13 libIREECompilerAggregateCAPI.dylib 0x000000010da4e1f8 mlir::PassInstrumentor::runAfterAnalysis(llvm::StringRef, mlir::TypeID, mlir::Operation*) + 4304
14 libIREECompilerAggregateCAPI.dylib 0x00000001111092f0 mlir::createInlinerPass(llvm::StringMap<mlir::OpPassManager, llvm::MallocAllocator>, std::__1::function<void (mlir::OpPassManager&)>) + 26604
15 libIREECompilerAggregateCAPI.dylib 0x0000000111109014 mlir::createInlinerPass(llvm::StringMap<mlir::OpPassManager, llvm::MallocAllocator>, std::__1::function<void (mlir::OpPassManager&)>) + 25872
16 libIREECompilerAggregateCAPI.dylib 0x0000000111103a90 mlir::createInlinerPass(llvm::StringMap<mlir::OpPassManager, llvm::MallocAllocator>, std::__1::function<void (mlir::OpPassManager&)>) + 3980
17 libIREECompilerAggregateCAPI.dylib 0x000000010da4a384 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) + 724
18 libIREECompilerAggregateCAPI.dylib 0x000000010da4ac3c mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) + 348
19 libIREECompilerAggregateCAPI.dylib 0x000000010da4c574 mlir::PassManager::run(mlir::Operation*) + 652
20 libIREECompilerAggregateCAPI.dylib 0x000000010d65a13c mlir::detail::walk(mlir::Operation*, llvm::function_ref<mlir::WalkResult (mlir::Operation*, mlir::WalkStage const&)>) + 2181736
21 libIREECompilerAggregateCAPI.dylib 0x000000010d659054 mlir::detail::walk(mlir::Operation*, llvm::function_ref<mlir::WalkResult (mlir::Operation*, mlir::WalkStage const&)>) + 2177408
22 dyld                               0x0000000100e0508c start + 520


Invoked with:
 iree-compile /Users/phawkins/.pyenv/versions/py310/lib/python3.10/site-packages/iree/compiler/tools/../_mlir_libs/iree-compile - --iree-input-type=mhlo --iree-vm-bytecode-module-output-format=flatbuffer-binary --iree-hal-target-backends=dylib --iree-mlir-to-vm-bytecode-module --iree-llvm-embedded-linker-path=/Users/phawkins/.pyenv/versions/py310/lib/python3.10/site-packages/iree/compiler/tools/../_mlir_libs/iree-lld --mlir-print-debuginfo --mlir-print-op-on-diagnostic=false

[hanhanW]

Why this is mhlo.pad op, not mhlo.slice op? The padding config is padding_config=((0, 0, 0), (-1, -1, 0))

[jpienaar]

Why this is mhlo.pad op, not mhlo.slice op? The padding config is padding_config=((0, 0, 0), (-1, -1, 0))

Let's file a bug for that canonicalization. And keep this one to the crash given this op.

[hawkinsp]

I don't have a strong opinion about whether this should be allowed by the MHLO spec, but right now it appears to be (in the sense that XLA accepts it if the IR is exported to XLA and that's as close as we have to a specification).

We should either clarify the MHLO spec and disallow this, or implement it.

And crashing isn't the right behavior either way...

[jpienaar]

The spec has been updated here and the above input seems to be passing for me, so issue potentially resolved.

(tested with

tools/iree-compile \
                                             --iree-input-type=mhlo \
                                             --iree-vm-bytecode-module-output-format=flatbuffer-binary \
                                             --iree-hal-target-backends=dylib \
                                             --iree-mlir-to-vm-bytecode-module \
                                             --mlir-print-debuginfo \
                                             --mlir-pass-pipeline-crash-reproducer=/tmp/iree/core-reproducer.mlir \
                                             9296.mlir -o /dev/null

and flatbuffer was generated)

[allieculp]

Assigning to Jacques officially as this is labeled as a P1 issue. It may be resolved?

[jpienaar]

Crash has been resolved.