llvm / llvm/llvm-project

[mlir] mlir-opt directly crashes at Block.cpp:250 with assertion `mightHaveTerminator()' failed.

Open
#205,239 1 comment 0 reactions 0 assignees View on GitHub
crash mlir
Dominant language
LLVM
Stars
40.5k
Forks
18.7k
PR merge metrics
PR metrics pending

Description

git version: c56e892f5c520b4474158837d7968e5c8dc3e4c1

MLIR Program:
```mlir
"builtin.module"() ({
"func.func"() <{function_type = () -> (), sym_name = "m"}> ({
%1 = "gpu.thread_id"() <{dimension = #gpu}> : () -> index
%2 = "gpu.warp_execute_on_lane_0"(%1) <{warp_size = 32 : i64}> ({
%15 = "arith.constant"() <{value = 1.000000e+00 : f32}> : () -> f32
%19 = "vector.from_elements"(%15, %15, %15, %15) {layout_result_0 = #xegpu.layout} : (f32, f32, f32, f32) -> vector<2x2xf32>
%20 = "vector.transpose"(%19) <{permutation = array}> {layout_operand_0 = #xegpu.layout, layout_result_0 = #xegpu.layout} : (vector<2x2xf32>) -> vector<2x2xf32>
}) {} : (index) -> vector<2x2xf32>
"func.return"() : () -> ()
}) : () -> ()
}) : () -> ()
```

Reproduce Command: `mlir-opt a.mlir`

StackTrace:
```
mlir-opt: /data2/dependency/dev/llvm-project/mlir/lib/IR/Block.cpp:250: mlir::Operation* mlir::Block::getTerminator(): Assertion `mightHaveTerminator()' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace and instructions to reproduce the bug.
Stack dump:
0. Program arguments: /data2/dependency/dev/llvm-project/build/bin/mlir-opt a.mlir
#0 0x000055ba98b3220f llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0x25cb20f)
#1 0x000055ba98b2e9dc llvm::sys::RunSignalHandlers() (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0x25c79dc)
#2 0x000055ba98b2eb77 SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
#3 0x00007f08a5f45420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420)
#4 0x00007f08a5a3000b raise (/lib/x86_64-linux-gnu/libc.so.6+0x4300b)
#5 0x00007f08a5a0f859 abort (/lib/x86_64-linux-gnu/libc.so.6+0x22859)
#6 0x00007f08a5a0f729 (/lib/x86_64-linux-gnu/libc.so.6+0x22729)
#7 0x00007f08a5a20fd6 (/lib/x86_64-linux-gnu/libc.so.6+0x33fd6)
#8 0x000055baa1bf28e5 (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0xb68b8e5)
#9 0x000055baa1bf3727 mlir::Block::getTerminator() (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0xb68c727)
#10 0x000055baa081a189 mlir::gpu::WarpExecuteOnLane0Op::verify() (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0xa2b3189)
#11 0x000055baa0853a60 mlir::Op::Impl, mlir::OpTrait::SingleBlock, mlir::OpTrait::SingleBlockImplicitTerminator::Impl, mlir::OpTrait::OpInvariants, mlir::BytecodeOpInterface::Trait, mlir::RegionBranchOpInterface::Trait, mlir::OpTrait::HasRecursiveMemoryEffects>::verifyInvariants(mlir::Operation*) (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0xa2eca60)
#12 0x000055baa08473e8 mlir::RegisteredOperationName::Model::verifyInvariants(mlir::Operation*) (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0xa2e03e8)
#13 0x000055baa1d06304 (anonymous namespace)::OperationVerifier::verifyOnEntrance(mlir::Operation&) Verifier.cpp:0:0
#14 0x000055baa1d08600 (anonymous namespace)::OperationVerifier::verifyOperation(mlir::Operation&) Verifier.cpp:0:0
#15 0x000055baa1d08f9f (anonymous namespace)::OperationVerifier::verifyOperation(mlir::Operation&) Verifier.cpp:0:0
#16 0x000055baa1d08f9f (anonymous namespace)::OperationVerifier::verifyOperation(mlir::Operation&) Verifier.cpp:0:0
#17 0x000055baa1d09f5f mlir::verify(mlir::Operation*, bool) (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0xb7a2f5f)
#18 0x000055baa15facef mlir::parseAsmSourceFile(llvm::SourceMgr const&, mlir::Block*, mlir::ParserConfig const&, mlir::AsmParserState*, mlir::AsmParserCodeCompleteContext*) (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0xb093cef)
#19 0x000055ba9fe4a255 mlir::parseSourceFile(std::shared_ptr const&, mlir::Block*, mlir::ParserConfig const&, mlir::LocationAttr*) (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0x98e3255)
#20 0x000055ba98bf6ecb performActions(llvm::raw_ostream&, std::shared_ptr const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#21 0x000055ba98bf8933 processBuffer(llvm::raw_ostream&, std::unique_ptr>, llvm::MemoryBufferRef, mlir::MlirOptMainConfig const&, mlir::DialectRegistry&, mlir::SourceMgrDiagnosticVerifierHandler*, llvm::ThreadPoolInterface*) MlirOptMain.cpp:0:0
#22 0x000055ba98bf8bbb llvm::LogicalResult llvm::function_ref>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>::callback_fn>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&)::'lambda'(std::unique_ptr>, llvm::MemoryBufferRef, llvm::raw_ostream&)>(long, std::unique_ptr>, llvm::MemoryBufferRef const&, llvm::raw_ostream&) MlirOptMain.cpp:0:0
#23 0x000055baa1d20cb5 mlir::splitAndProcessBuffer(std::unique_ptr>, llvm::function_ref>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>, llvm::raw_ostream&, llvm::StringRef, llvm::StringRef) (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0xb7b9cb5)
#24 0x000055ba98beecc4 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (.part.0) MlirOptMain.cpp:0:0
#25 0x000055ba98bf90d7 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0x26920d7)
#26 0x000055ba98bf9313 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0x2692313)
#27 0x000055ba98a5eb8c main (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0x24f7b8c)
#28 0x00007f08a5a11083 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x24083)
#29 0x000055ba98b0e91e _start (/data2/dependency/dev/llvm-project/build/bin/mlir-opt+0x25a791e)
Aborted (core dumped)
```

Contributor guide

Open the contributing guide

Research direction

Reproduce the crash with mlir-opt a.mlir using the supplied MLIR program. Start with mlir/lib/IR/Block.cpp at getTerminator() and gpu::WarpExecuteOnLane0Op::verify(), the entry points shown in the stack trace, then trace the verifier path. Done means the reproducer no longer aborts and has a regression test.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
compilers
Issue type
Bug
Difficulty
3/5
Estimated time
1-2 days
Activity status
Quiet
Clarity
Mostly clear
Newbie friendliness
55/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.