| Issue |
128837
|
| Summary |
[MLIR] --gpu-kernel-outlining triggers Assertion `isa<To>(Val) && "cast<Ty>() argument of incompatible type!"' failed.
|
| Labels |
mlir
|
| Assignees |
|
| Reporter |
RiRi114
|
Test on commit : 0f9720a6
steps to reproduce:
`
mlir-opt temp.mlir --gpu-kernel-outlining
`
test case:
```mlir
module {
func.func @kernel_function() {
// Placeholder for kernel functionality, operates on input data
%data = "" : memref<4x8xf32>
%result = memref.alloc() : memref<4xf32>
%c1 = arith.constant 1 : index
gpu.launch blocks(%bx, %by, %bz) in (%grid_x = %c1, %grid_y = %c1, %grid_z = %c1)
threads(%tx, %ty, %tz) in (%block_x = %c1, %block_y = %c1, %block_z = %c1) {
%val = memref.load %data[%bx, %tx] : memref<4x8xf32>
%sum_result = gpu.all_reduce add %val uniform {} : (f32) -> (f32)
memref.store %sum_result, %result[%bx] : memref<4xf32>
gpu.terminator
}
return
}
func.func @main() {
%data_init = memref.alloc() : memref<4x8xf32>
%result_init = memref.alloc() : memref<4xf32>
// Kernel launch with specific grid and block sizes
%c2 = arith.constant 2 : index
gpu.launch blocks(%bx, %by, %bz) in (%grid_x = %c2, %grid_y = %c2, %grid_z = %c2)
threads(%tx, %ty, %tz) in (%block_x = %c2, %block_y = %c2, %block_z = %c2) {
// Kernel computation
func.call @kernel_function() : () -> ()
gpu.terminator
}
return
}
}
```
crash trace:
```
mlir-opt: /home/fuzzing/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: /home/fuzzing/llvm-project/build/bin/mlir-opt temp.mlir --gpu-kernel-outlining
#0 0x000000000113ce77 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x113ce77)
#1 0x000000000113aa0e llvm::sys::RunSignalHandlers() (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x113aa0e)
#2 0x000000000113d855 SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
#3 0x0000718280502520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
#4 0x00007182805569fc pthread_kill (/lib/x86_64-linux-gnu/libc.so.6+0x969fc)
#5 0x0000718280502476 gsignal (/lib/x86_64-linux-gnu/libc.so.6+0x42476)
#6 0x00007182804e87f3 abort (/lib/x86_64-linux-gnu/libc.so.6+0x287f3)
#7 0x00007182804e871b (/lib/x86_64-linux-gnu/libc.so.6+0x2871b)
#8 0x00007182804f9e96 (/lib/x86_64-linux-gnu/libc.so.6+0x39e96)
#9 0x000000000193c5be 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 0x00000000012714f7 mlir::WalkResult mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<mlir::WalkResult (mlir::Operation*)>, mlir::WalkOrder) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x12714f7)
#11 0x000000000193ae42 (anonymous namespace)::GpuKernelOutliningPass::runOnOperation() KernelOutlining.cpp:0:0
#12 0x00000000042452f7 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x42452f7)
#13 0x0000000004245b61 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x4245b61)
#14 0x000000000424825b mlir::PassManager::run(mlir::Operation*) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x424825b)
#15 0x00000000042408af performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#16 0x0000000004240503 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&)::$_3>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer> >, llvm::raw_ostream&) MlirOptMain.cpp:0:0
#17 0x00000000042ec4e8 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/fuzzing/llvm-project/build/bin/mlir-opt+0x42ec4e8)
#18 0x000000000423a151 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer> >, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x423a151)
#19 0x000000000423a403 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x423a403)
#20 0x000000000423a612 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x423a612)
#21 0x000000000111a8d7 main (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x111a8d7)
#22 0x00007182804e9d90 (/lib/x86_64-linux-gnu/libc.so.6+0x29d90)
#23 0x00007182804e9e40 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x29e40)
#24 0x000000000111a435 _start (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x111a435)
Aborted (core dumped)
```
![]()
_______________________________________________
llvm-bugs mailing list
llvm-bugs@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
