| Issue |
149318
|
| Summary |
[MLIR] crashed in `-gpu-kernel-outlining` pass with error message: Assertion `isa<To>(Val) && "cast<Ty>() argument of incompatible type!"' failed
|
| Labels |
mlir
|
| Assignees |
|
| Reporter |
sweead
|
test commit: [66da9f38f](https://github.com/llvm/llvm-project/commit/66da9f38f374e786b2f1c0ecdab0b651c94c4f27)
Step to reproduce:
```
mlir-opt test.mlir -gpu-kernel-outlining
```
test case:
```
#map = affine_map<(d0)[s0, s1] -> ((d0 - s0) ceildiv s1)>
module {
func.func @main(%arg0: memref<89xi32, strided<[?], offset: ?>>, %arg1: memref<1xi32, strided<[?], offset: ?>>, %arg2: memref<59x3x9x39xf32, strided<[?, ?, ?, ?], offset: ?>>, %arg3: memref<1x1x9x1xf32, strided<[?, ?, ?, ?], offset: ?>>, %arg4: memref<i8, strided<[], offset: ?>>, %arg5: memref<i8, strided<[], offset: ?>>){
%c3 = arith.constant 3 : index
%c59 = arith.constant 59 : index
%c39 = arith.constant 39 : index
%c9 = arith.constant 9 : index
%c1 = arith.constant 1 : index
%c0 = arith.constant 0 : index
gpu.launch blocks(%arg6, %arg7, %arg8) in (%arg12 = %c59, %arg13 = %c1, %arg14 = %c1) threads(%arg9, %arg10, %arg11) in (%arg15 = %c3, %arg16 = %c1, %arg17 = %c1) {
scf.for %arg18 = %c0 to %c9 step %c1 {
%c1_8 = arith.constant 1 : index
%8 = affine.apply #map(%c39)[%c0, %c1]
gpu.launch blocks(%arg19, %arg20, %arg21) in (%arg25 = %8, %arg26 = %c1_8, %arg27 = %c1_8) threads(%arg22, %arg23, %arg24) in (%arg28 = %c1_8, %arg29 = %c1_8, %arg30 = %c1_8) {
gpu.terminator
} {SCFToGPU_visited}
}
gpu.terminator
} {SCFToGPU_visited}
return
}
}
```
Crash backtrace:
```
mlir-opt: /home/workdir/llvm-project/llvm/include/llvm/Support/Casting.h:566: decltype(auto) llvm::cast(const From &) [To = mlir::FlatSymbolRefAttr, From = mlir::SymbolRefAttr]: Assertion `isa<To>(Val) && "cast<Ty>() argument of incompatible type!"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0. Program arguments: mlir-opt test.mlir -gpu-kernel-outlining
#0 0x00005620458c1138 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/workdir/llvm-project/build/bin/mlir-opt+0x18bf138)
#1 0x00005620458be7b5 llvm::sys::RunSignalHandlers() (/home/workdir/llvm-project/build/bin/mlir-opt+0x18bc7b5)
#2 0x00005620458c2251 SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
#3 0x00007fa8d0b38520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
#4 0x00007fa8d0b8c9fc pthread_kill (/lib/x86_64-linux-gnu/libc.so.6+0x969fc)
#5 0x00007fa8d0b38476 gsignal (/lib/x86_64-linux-gnu/libc.so.6+0x42476)
#6 0x00007fa8d0b1e7f3 abort (/lib/x86_64-linux-gnu/libc.so.6+0x287f3)
#7 0x00007fa8d0b1e71b (/lib/x86_64-linux-gnu/libc.so.6+0x2871b)
#8 0x00007fa8d0b2fe96 (/lib/x86_64-linux-gnu/libc.so.6+0x39e96)
#9 0x000056204621a54d mlir::WalkResult llvm::function_ref<mlir::WalkResult (mlir::Operation*)>::callback_fn<std::enable_if<!llvm::is_one_of<mlir::gpu::LaunchOp, mlir::Operation*, mlir::Region*, mlir::Block*>::value && std::is_same<mlir::WalkResult, mlir::WalkResult>::value, mlir::WalkResult>::type mlir::detail::walk<(mlir::WalkOrder)1, mlir::ForwardIterator, (anonymous namespace)::GpuKernelOutliningPass::runOnOperation()::'lambda'(mlir::gpu::LaunchOp), mlir::gpu::LaunchOp, mlir::WalkResult>(mlir::Operation*, (anonymous namespace)::GpuKernelOutliningPass::runOnOperation()::'lambda'(mlir::gpu::LaunchOp)&&)::'lambda'(mlir::Operation*)>(long, mlir::Operation*) KernelOutlining.cpp:0:0
#10 0x0000562045a01e57 mlir::WalkResult mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<mlir::WalkResult (mlir::Operation*)>, mlir::WalkOrder) (/home/workdir/llvm-project/build/bin/mlir-opt+0x19ffe57)
#11 0x0000562046219122 (anonymous namespace)::GpuKernelOutliningPass::runOnOperation() KernelOutlining.cpp:0:0
#12 0x000056204915e1a3 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (/home/workdir/llvm-project/build/bin/mlir-opt+0x515c1a3)
#13 0x000056204915ea42 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+0x515ca42)
#14 0x000056204916121e mlir::PassManager::run(mlir::Operation*) (/home/workdir/llvm-project/build/bin/mlir-opt+0x515f21e)
#15 0x000056204915918b performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#16 0x0000562049158de4 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
#17 0x0000562049208ce5 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+0x5206ce5)
#18 0x0000562049151fd2 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+0x514ffd2)
#19 0x0000562049152288 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (/home/workdir/llvm-project/build/bin/mlir-opt+0x5150288)
#20 0x00005620491524a2 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (/home/workdir/llvm-project/build/bin/mlir-opt+0x51504a2)
#21 0x000056204589fb07 main (/home/workdir/llvm-project/build/bin/mlir-opt+0x189db07)
#22 0x00007fa8d0b1fd90 (/lib/x86_64-linux-gnu/libc.so.6+0x29d90)
#23 0x00007fa8d0b1fe40 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x29e40)
#24 0x000056204589f5a5 _start (/home/workdir/llvm-project/build/bin/mlir-opt+0x189d5a5)
Aborted (core dumped)
```
![]()
_______________________________________________
llvm-bugs mailing list
llvm-bugs@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
