| Issue |
131132
|
| Summary |
[MLIR] crashed on gpu::LaunchOp when using '-test-ir-visitors' pass
|
| Labels |
mlir
|
| Assignees |
|
| Reporter |
sweead
|
Test commit: https://github.com/llvm/llvm-project/commit/9e91725fd4d4ee30e98ab2682f93b423590a4ade
Steps to reproduce:
```shell
mlir-opt test.mlir -test-ir-visitors
```
Test case:
```mlir
module {
func.func @main(%arg0: tensor<4x4xf32>, %arg1: tensor<*xf32>, %arg2: tensor<*xf32>, %arg3: tensor<2x3xi32>, %arg4: tensor<2x3xi32>) attributes {pattern_driver_all_erased = true, pattern_driver_changed = false} {
%c4 = arith.constant 4 : index
%c0 = arith.constant 0 : index
%0 = llvm.mlir.constant(0 : i32) : i32
%1 = llvm.mlir.poison : vector<128xi32>
%2 = llvm.mlir.constant(1 : index) : i64
%3 = llvm.mlir.constant(dense<3.40282347E+38> : vector<128xf32>) : vector<128xf32>
%4 = bufferization.to_memref %arg0 : tensor<4x4xf32> to memref<4x4xf32, strided<[?, ?], offset: ?>, #spirv.storage_class<StorageBuffer>>
%alloc = memref.alloc() {alignment = 64 : i64} : memref<4xf32, #spirv.storage_class<StorageBuffer>>
%5 = builtin.unrealized_conversion_cast %alloc : memref<4xf32, #spirv.storage_class<StorageBuffer>> to !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%6 = builtin.unrealized_conversion_cast %2 : i64 to index
gpu.launch blocks(%arg5, %arg6, %arg7) in (%arg11 = %6, %arg12 = %6, %arg13 = %6) threads(%arg8, %arg9, %arg10) in (%arg14 = %6, %arg15 = %6, %arg16 = %6) {
%7 = llvm.trunc %2 : i64 to i32
%8 = llvm.insertelement %7, %1[%0 : i32] : vector<128xi32>
%9 = llvm.shufflevector %8, %1 [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0] : vector<128xi32>
%10 = llvm.icmp "sgt" %9, %1 : vector<128xi32>
%11 = llvm.extractvalue %5[1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%12 = llvm.getelementptr %11[%2] : (!llvm.ptr, i64) -> !llvm.ptr, f32
llvm.intr.masked.store %3, %12, %10 {alignment = 4 : i32} : vector<128xf32>, vector<128xi1> into !llvm.ptr
gpu.terminator
}
return
}
}
```
Crash trace:
```console
Block forward dominance post-order visits
Visiting block ^bb0mlir-opt: /home/workdir/llvm-project/llvm/include/llvm/ADT/ArrayRef.h:84: llvm::ArrayRef<mlir::BlockArgument>::ArrayRef(const T *, const T *) [T = mlir::BlockArgument]: Assertion `begin <= end' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0. Program arguments: /home/workdir/llvm-project/build/bin/./mlir-opt test.mlir -test-ir-visitors
#0 0x000055a5ee45b228 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x12d1228)
#1 0x000055a5ee458d4e llvm::sys::RunSignalHandlers() (/home/workdir/llvm-project/build/bin/./mlir-opt+0x12ced4e)
#2 0x000055a5ee45bc31 SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
#3 0x00007fb094c7f520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
#4 0x00007fb094cd39fc pthread_kill (/lib/x86_64-linux-gnu/libc.so.6+0x969fc)
#5 0x00007fb094c7f476 gsignal (/lib/x86_64-linux-gnu/libc.so.6+0x42476)
#6 0x00007fb094c657f3 abort (/lib/x86_64-linux-gnu/libc.so.6+0x287f3)
#7 0x00007fb094c6571b (/lib/x86_64-linux-gnu/libc.so.6+0x2871b)
#8 0x00007fb094c76e96 (/lib/x86_64-linux-gnu/libc.so.6+0x39e96)
#9 0x000055a5eeb49ebf mlir::gpu::LaunchOp::getPrivateAttributions() (/home/workdir/llvm-project/build/bin/./mlir-opt+0x19bfebf)
#10 0x000055a5eeb49473 mlir::gpu::LaunchOp::verifyRegions() (/home/workdir/llvm-project/build/bin/./mlir-opt+0x19bf473)
#11 0x000055a5eec19055 mlir::Op<mlir::gpu::LaunchOp, mlir::OpTrait::OneRegion, mlir::OpTrait::VariadicResults, mlir::OpTrait::ZeroSuccessors, mlir::OpTrait::AtLeastNOperands<6u>::Impl, mlir::OpTrait::AttrSizedOperandSegments, mlir::OpTrait::OpInvariants, mlir::BytecodeOpInterface::Trait, mlir::OpTrait::AffineScope, mlir::OpTrait::AutomaticAllocationScope, mlir::InferIntRangeInterface::Trait, mlir::gpu::AsyncOpInterface::Trait, mlir::OpTrait::HasRecursiveMemoryEffects>::verifyRegionInvariants(mlir::Operation*) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x1a8f055)
#12 0x000055a5eec174f0 mlir::RegisteredOperationName::Model<mlir::gpu::LaunchOp>::verifyRegionInvariants(mlir::Operation*) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x1a8d4f0)
#13 0x000055a5f171439b (anonymous namespace)::OperationVerifier::verifyOpAndDominance(mlir::Operation&) Verifier.cpp:0:0
#14 0x000055a5f17137b6 mlir::verify(mlir::Operation*, bool) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x45897b6)
#15 0x000055a5f160a333 mlir::AsmState::AsmState(mlir::Operation*, mlir::OpPrintingFlags const&, llvm::DenseMap<mlir::Operation*, std::pair<unsigned int, unsigned int>, llvm::DenseMapInfo<mlir::Operation*, void>, llvm::detail::DenseMapPair<mlir::Operation*, std::pair<unsigned int, unsigned int>>>*, mlir::FallbackAsmResourceMap*) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x4480333)
#16 0x000055a5f1616a45 mlir::Block::printAsOperand(llvm::raw_ostream&, bool) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x448ca45)
#17 0x000055a5f1be2383 printBlock(mlir::Block*) TestVisitors.cpp:0:0
#18 0x000055a5f1be6f0c void llvm::function_ref<void (mlir::Block*)>::callback_fn<testNoSkipErasureCallbacks(mlir::Operation*)::$_1>(long, mlir::Block*) TestVisitors.cpp:0:0
#19 0x000055a5f1134f9e void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Block*)>, mlir::WalkOrder) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x3faaf9e)
#20 0x000055a5f1134f83 void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Block*)>, mlir::WalkOrder) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x3faaf83)
#21 0x000055a5f1be2086 (anonymous namespace)::TestIRVisitorsPass::runOnOperation() TestVisitors.cpp:0:0
#22 0x000055a5f1551cd3 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x43c7cd3)
#23 0x000055a5f1552572 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x43c8572)
#24 0x000055a5f1554d4e mlir::PassManager::run(mlir::Operation*) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x43cad4e)
#25 0x000055a5f154d2cb performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#26 0x000055a5f154cf23 llvm::LogicalResult llvm::function_ref<llvm::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, 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::raw_ostream&) MlirOptMain.cpp:0:0
#27 0x000055a5f15f8475 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::raw_ostream&)>, llvm::raw_ostream&, llvm::StringRef, llvm::StringRef) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x446e475)
#28 0x000055a5f1546b82 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x43bcb82)
#29 0x000055a5f1546e33 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x43bce33)
#30 0x000055a5f1547042 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x43bd042)
#31 0x000055a5ee43a9af main (/home/workdir/llvm-project/build/bin/./mlir-opt+0x12b09af)
#32 0x00007fb094c66d90 (/lib/x86_64-linux-gnu/libc.so.6+0x29d90)
#33 0x00007fb094c66e40 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x29e40)
#34 0x000055a5ee43a505 _start (/home/workdir/llvm-project/build/bin/./mlir-opt+0x12b0505)
```
![]()
_______________________________________________
llvm-bugs mailing list
llvm-bugs@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
