| Issue |
171303
|
| Summary |
[MLIR] Crash triggered by MLIR lowering pass
|
| Labels |
mlir
|
| Assignees |
|
| Reporter |
clearlove2077
|
Hi, I found a crash bug when lowering programs with MLIR. Here is a code example that reproduces the issue:
**Input file name: input.mlir**
```llvm
module {
gpu.module @vector_add_kernel {
gpu.func @vector_add_local_tile_multi_elem_per_thread_half(%z: memref<?xf16>, %num: i32, %x: memref<?xf16>, %y: memref<?xf16>, %a: f16, %b: f16, %c: f16) {
%kNumElemPerThread = arith.constant 8 : i32
%blockDim_x = gpu.global_id x
%threadIdx_x = gpu.thread_id x
%blockIdx_x = gpu.block_id x
%blockDim_x_i32 = arith.index_cast %blockDim_x : index to i32
%threadIdx_x_i32 = arith.index_cast %threadIdx_x : index to i32
%blockIdx_x_i32 = arith.index_cast %blockIdx_x : index to i32
%grid_index = arith.muli %blockIdx_x_i32, %blockDim_x_i32 : i32
%idx = arith.addi %threadIdx_x_i32, %grid_index : i32
%num_div_elems = arith.divsi %num, %kNumElemPerThread : i32
%cond = arith.cmpi slt, %idx, %num_div_elems : i32
scf.if %cond {
%offset = arith.muli %idx, %kNumElemPerThread : i32
%offset_index = arith.index_cast %offset : i32 to index
%tx_data = vector.load %x[%offset_index] : memref<?xf16>, vector<8xf16>
%ty_data = vector.load %y[%offset_index] : memref<?xf16>, vector<8xf16>
%a_extended = arith.extf %a : f16 to f32
%b_extended = arith.extf %b : f16 to f32
%c_extended = arith.extf %c : f16 to f32
%a_broadcast = vector.broadcast %a_extended : f32 to vector<8xf32>
%b_broadcast = vector.broadcast %b_extended : f32 to vector<8xf32>
%c_broadcast = vector.broadcast %c_extended : f32 to vector<8xf32>
%a_truncated = arith.truncf %a_broadcast : vector<8xf32> to vector<8xf16>
%b_truncated = arith.truncf %b_broadcast : vector<8xf32> to vector<8xf16>
%c_truncated = arith.truncf %c_broadcast : vector<8xf32> to vector<8xf16>
%product_a_x = arith.mulf %tx_data, %a_truncated : vector<8xf16>
%product_b_y = arith.mulf %ty_data, %b_truncated : vector<8xf16>
%sum_bc = arith.addf %product_b_y, %c_truncated : vector<8xf16>
%final_result = arith.addf %product_a_x, %sum_bc : vector<8xf16>
vector.store %final_result, %z[%offset_index] : memref<?xf16>, vector<8xf16>
}
gpu.return
}
}
}
```
I encountered a crash bug when running MLIR with the following command.
### Used Command
```bash
mlir-opt ./input.mlir -gpu-kernel-outlining --convert-scf-to-cf --convert-gpu-to-nvvm -reconcile-unrealized-casts --nvvm-attach-target=chip=sm_75 --gpu-module-to-binary -gpu-to-llvm --convert-cf-to-llvm --convert-index-to-llvm --convert-arith-to-llvm --convert-math-to-llvm --convert-func-to-llvm --finalize-memref-to-llvm -convert-nvvm-to-llvm -reconcile-unrealized-casts
```
### Crash Log
```bash
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: mlir-opt ./input.mlir -gpu-kernel-outlining --convert-scf-to-cf --convert-gpu-to-nvvm -reconcile-unrealized-casts --nvvm-attach-target=chip=sm_75 --gpu-module-to-binary -gpu-to-llvm --convert-cf-to-llvm --convert-index-to-llvm --convert-arith-to-llvm --convert-math-to-llvm --convert-func-to-llvm --finalize-memref-to-llvm -convert-nvvm-to-llvm -allow-unregistered-dialect -reconcile-unrealized-casts
#0 0x000055d0a911a508 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x1b3f508)
#1 0x000055d0a9117c35 llvm::sys::RunSignalHandlers() (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x1b3cc35)
#2 0x000055d0a911b611 SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
#3 0x00007f9a53c45330 (/lib/x86_64-linux-gnu/libc.so.6+0x45330)
#4 0x00007f9a53c9eb2c __pthread_kill_implementation ./nptl/pthread_kill.c:44:76
#5 0x00007f9a53c9eb2c __pthread_kill_internal ./nptl/pthread_kill.c:78:10
#6 0x00007f9a53c9eb2c pthread_kill ./nptl/pthread_kill.c:89:10
#7 0x00007f9a53c4527e raise ./signal/../sysdeps/posix/raise.c:27:6
#8 0x00007f9a53c288ff abort ./stdlib/abort.c:81:7
#9 0x000055d0a910263d llvm::report_fatal_error(llvm::Twine const&, bool) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x1b2763d)
#10 0x000055d0ae46c62c (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x6e9162c)
#11 0x000055d0ae45c879 mlir::index::MulOp::create(mlir::OpBuilder&, mlir::Location, mlir::Value, mlir::Value) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x6e81879)
#12 0x000055d0aaac678b (anonymous namespace)::GpuGlobalIdRewriter::matchAndRewrite(mlir::gpu::GlobalIdOp, mlir::PatternRewriter&) const GlobalIdRewriter.cpp:0:0
#13 0x000055d0b032e39e void llvm::function_ref<void ()>::callback_fn<mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<llvm::LogicalResult (mlir::Pattern const&)>)::$_0>(long) PatternApplicator.cpp:0:0
#14 0x000055d0b032a8eb mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<llvm::LogicalResult (mlir::Pattern const&)>) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x8d4f8eb)
#15 0x000055d0b03065bb (anonymous namespace)::GreedyPatternRewriteDriver::processWorklist() GreedyPatternRewriteDriver.cpp:0:0
#16 0x000055d0b0303150 mlir::applyPatternsGreedily(mlir::Region&, mlir::FrozenRewritePatternSet const&, mlir::GreedyRewriteConfig, bool*) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x8d28150)
#17 0x000055d0aa61ced8 (anonymous namespace)::LowerGpuOpsToNVVMOpsPass::runOnOperation() LowerGpuOpsToNVVMOps.cpp:0:0
#18 0x000055d0b0389bc1 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x8daebc1)
#19 0x000055d0b038ac07 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x8dafc07)
#20 0x000055d0b0395dae auto void mlir::parallelForEach<__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_12>(mlir::MLIRContext*, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_12&&)::'lambda'(__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>&&)::operator()<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo&>(__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>&&) const Pass.cpp:0:0
#21 0x000055d0b038d8eb mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x8db28eb)
#22 0x000055d0b038a01f mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x8daf01f)
#23 0x000055d0b038ac07 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x8dafc07)
#24 0x000055d0b03921a3 mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x8db71a3)
#25 0x000055d0b03917b2 mlir::PassManager::run(mlir::Operation*) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x8db67b2)
#26 0x000055d0a91bf603 performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#27 0x000055d0a91be864 llvm::LogicalResult llvm::function_ref<llvm::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>::callback_fn<mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&)::$_0>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef const&, llvm::raw_ostream&) MlirOptMain.cpp:0:0
#28 0x000055d0b06c6928 mlir::splitAndProcessBuffer(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<llvm::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>, llvm::raw_ostream&, llvm::StringRef, llvm::StringRef) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x90eb928)
#29 0x000055d0a91b3dac mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x1bd8dac)
#30 0x000055d0a91b40e6 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x1bd90e6)
#31 0x000055d0a91b42f2 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x1bd92f2)
#32 0x000055d0a910224f main (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x1b2724f)
#33 0x00007f9a53c2a1ca __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#34 0x00007f9a53c2a28b call_init ./csu/../csu/libc-start.c:128:20
#35 0x00007f9a53c2a28b __libc_start_main ./csu/../csu/libc-start.c:347:5
#36 0x000055d0a9101d75 _start (/mnt/raid/home/username/bin/mlir/bin/mlir-opt+0x1b26d75)
Aborted
```
My MLIR version:
```bash
mlir-opt --version
LLVM (http://llvm.org/):
LLVM version 22.0.0git
Optimized build with assertions.
```
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
