| Issue |
185357
|
| Summary |
[MLIR] Crash during GPU kernel outlining when cloning operation
|
| Labels |
mlir
|
| Assignees |
|
| Reporter |
Zustin
|
Hi, I found a crash bug when lowering an MLIR program with the Vulkan pipeline. See the following code example:
Git version: da8d1b52d0d344068a8bd3e46a8af3a692d2efc2
```llvm
module attributes {gpu.container_module} {
gpu.module @kernels {
gpu.func @increment_kernel(%arg0: memref<?xi32>, %arg1: i32, %arg2: memref<?xi1>) kernel {
%c0 = arith.constant 0 : i32
%tid = gpu.thread_id x
%bid = gpu.block_id x
%bdim = gpu.block_dim x
%tid_i32 = arith.index_cast %tid : index to i32
%bid_i32 = arith.index_cast %bid : index to i32
%bdim_i32 = arith.index_cast %bdim : index to i32
%idx = arith.addi %bid, %tid : index
%val = memref.load %arg0[%idx] : memref<?xi32>
%cond = arith.cmpi slt, %val, %arg1 : i32
memref.store %cond, %arg2[%idx] : memref<?xi1>
%new_val = arith.addi %val, %arg1 : i32
memref.store %new_val, %arg0[%idx] : memref<?xi32>
gpu.return
}
}
func.func @main() {
%zero = arith.constant 0 : i32
%blocksz = arith.constant 32 : i32
%_one_ = arith.constant 1 : i32
%total = arith.constant 131072 : i32
%threshold = arith.constant 50 : i32
%h_data = memref.alloc() : memref<131072xi32>
%d_data = gpu.alloc() : memref<131072xi32>
%h_flag = memref.alloc() : memref<131072xi1>
%d_flag = gpu.alloc() : memref<131072xi1>
%grid = arith.divsi %total, %blocksz : i32
%grid_mod = arith.remsi %grid, %threshold : i32
%st = gpu.wait async []
%memcpy_token0 = gpu.memcpy async [%st] %d_data, %h_data : memref<131072xi32>, memref<131072xi32>
%memcpy_token1 = gpu.memcpy async [%st] %d_flag, %h_flag : memref<131072xi1>, memref<131072xi1>
%grid_mod_index = arith.index_cast %grid_mod : i32 to index
%blocksz_index = arith.index_cast %blocksz : i32 to index
%_one_index_ = arith.index_cast %one : i32 to index
%total_index = arith.index_cast %total : i32 to index
%launch_token = gpu.launch async [%memcpy_token1]
blocks(%bX, %bY, %bZ) in (%g1 = %grid_mod_index, %g2 = %one_index, %g3 = %one_index)
threads(%tX, %tY, %tZ) in (%b1 = %blocksz_index, %b2 = %one_index, %b3 = %one_index) {
gpu.launch_func @kernels::@increment_kernel
blocks in (%bX, %bY, %bZ)
threads in (%tX, %tY, %tZ)
args(%d_data : memref<131072xi32>, %threshold : i32, %d_flag : memref<131072xi1>)
gpu.terminator
}
%memcpy_token2 = gpu.memcpy async [%launch_token] %h_data, %d_data : memref<131072xi32>, memref<131072xi32>
%memcpy_token3 = gpu.memcpy async [%launch_token] %h_flag, %d_flag : memref<131072xi1>, memref<131072xi1>
%init_val = arith.constant 0 : i32
%zero_index = arith.index_cast %zero : i32 to index
%_one_index_2_ = arith.index_cast %one : i32 to index
%loop_result = scf.for %i = %zero_index to %total_index step %one_index_2 iter_args(%acc = %init_val) -> i32 {
%wait_token = gpu.wait async [%memcpy_token2]
%new_acc = arith.addi %acc, %one : i32
scf.yield %new_acc : i32
}
%final_wait = gpu.wait async [%memcpy_token3]
memref.dealloc %h_data : memref<131072xi32>
gpu.dealloc %d_data : memref<131072xi32>
memref.dealloc %h_flag : memref<131072xi1>
gpu.dealloc %d_flag : memref<131072xi1>
return
}
}
```
#### Stack Trace
```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 asyncAPI.mlir -test-vulkan-runner-pipeline
#0 0x000056c677f30dbd llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) /llvm/lib/Support/Unix/Signals.inc:880:11
#1 0x000056c677f312eb PrintStackTraceSignalHandler(void*) /llvm/lib/Support/Unix/Signals.inc:962:1
#2 0x000056c677f2f184 llvm::sys::RunSignalHandlers() /llvm/lib/Support/Signals.cpp:108:5
#3 0x000056c677f31a89 SignalHandler(int, siginfo_t*, void*) /llvm/lib/Support/Unix/Signals.inc:448:38
#4 0x000070a08a445330 (/lib/x86_64-linux-gnu/libc.so.6+0x45330)
#5 0x000056c67813eee4 mlir::Operation::getNumOperands() /mlir/include/mlir/IR/Operation.h:376:24
#6 0x000056c686dc7306 mlir::Operation::clone(mlir::IRMapping&, mlir::Operation::CloneOptions const&) /mlir/lib/IR/Operation.cpp:731:22
#7 0x000056c686dc7825 mlir::Operation::clone(mlir::Operation::CloneOptions const&) /mlir/lib/IR/Operation.cpp:763:10
#8 0x000056c67b20e2a6 (anonymous namespace)::GpuKernelOutliningPass::createKernelModule(mlir::gpu::LaunchOp, mlir::gpu::GPUFuncOp, mlir::SymbolTable const&) /mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp:443:53
#9 0x000056c67b20dd4c (anonymous namespace)::GpuKernelOutliningPass::runOnOperation()::'lambda'(mlir::gpu::LaunchOp)::operator()(mlir::gpu::LaunchOp) const /mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp:378:29
#10 0x000056c67b20db7c 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*)::operator()(mlir::Operation*) const /mlir/include/mlir/IR/Visitors.h:341:14
#11 0x000056c67b20db1d 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*) /llvm/include/llvm/ADT/STLFunctionalExtras.h:46:12
#12 0x000056c6780e2a81 llvm::function_ref<mlir::WalkResult (mlir::Operation*)>::operator()(mlir::Operation*) const /llvm/include/llvm/ADT/STLFunctionalExtras.h:69:12
#13 0x000056c6780e2a41 mlir::WalkResult mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<mlir::WalkResult (mlir::Operation*)>, mlir::WalkOrder) /mlir/include/mlir/IR/Visitors.h:246:12
#14 0x000056c6780e29d6 mlir::WalkResult mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<mlir::WalkResult (mlir::Operation*)>, mlir::WalkOrder) /mlir/include/mlir/IR/Visitors.h:239:13
#15 0x000056c67b20daba 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)&&) /mlir/include/mlir/IR/Visitors.h:344:10
#16 0x000056c67b20da5d std::enable_if<llvm::function_traits<std::decay<(anonymous namespace)::GpuKernelOutliningPass::runOnOperation()::'lambda'(mlir::gpu::LaunchOp)>::type>::num_args == 1, mlir::WalkResult>::type mlir::Operation::walk<(mlir::WalkOrder)1, mlir::ForwardIterator, (anonymous namespace)::GpuKernelOutliningPass::runOnOperation()::'lambda'(mlir::gpu::LaunchOp), mlir::WalkResult>((anonymous namespace)::GpuKernelOutliningPass::runOnOperation()::'lambda'(mlir::gpu::LaunchOp)&&) /mlir/include/mlir/IR/Operation.h:827:12
#17 0x000056c67b20da30 std::enable_if<llvm::function_traits<std::decay<(anonymous namespace)::GpuKernelOutliningPass::runOnOperation()::'lambda'(mlir::gpu::LaunchOp)>::type>::num_args == 1, mlir::WalkResult>::type mlir::OpState::walk<(mlir::WalkOrder)1, mlir::ForwardIterator, (anonymous namespace)::GpuKernelOutliningPass::runOnOperation()::'lambda'(mlir::gpu::LaunchOp), mlir::WalkResult>((anonymous namespace)::GpuKernelOutliningPass::runOnOperation()::'lambda'(mlir::gpu::LaunchOp)&&) /mlir/include/mlir/IR/OpDefinition.h:170:19
#18 0x000056c67b20d5eb (anonymous namespace)::GpuKernelOutliningPass::runOnOperation() /mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp:360:34
#19 0x000056c6868d9fa4 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_3::operator()() const /mlir/lib/Pass/Pass.cpp:0:19
#20 0x000056c6868d9f45 void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_3>(long) /llvm/include/llvm/ADT/STLFunctionalExtras.h:46:5
#21 0x000056c677f57189 llvm::function_ref<void ()>::operator()() const /llvm/include/llvm/ADT/STLFunctionalExtras.h:69:5
#22 0x000056c6868dd44b void mlir::MLIRContext::executeAction<mlir::PassExecutionAction, mlir::Pass&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pass&) /mlir/include/mlir/IR/MLIRContext.h:291:3
#23 0x000056c6868d13ee mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /mlir/lib/Pass/Pass.cpp:619:17
#24 0x000056c6868d1b0a mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /mlir/lib/Pass/Pass.cpp:688:16
#25 0x000056c6868d73ed mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) /mlir/lib/Pass/Pass.cpp:1123:10
#26 0x000056c6868d6976 mlir::PassManager::run(mlir::Operation*) /mlir/lib/Pass/Pass.cpp:1097:60
#27 0x000056c67803be7b performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) /mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:596:17
#28 0x000056c67803b551 processBuffer(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef, mlir::MlirOptMainConfig const&, mlir::DialectRegistry&, mlir::SourceMgrDiagnosticVerifierHandler*, llvm::ThreadPoolInterface*) /mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:681:12
#29 0x000056c67803b20c mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&)::$_0::operator()(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef, llvm::raw_ostream&) const /mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:779:12
#30 0x000056c67803b116 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&) /llvm/include/llvm/ADT/STLFunctionalExtras.h:46:12
#31 0x000056c686e436fa llvm::function_ref<llvm::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>::operator()(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef const&, llvm::raw_ostream&) const /llvm/include/llvm/ADT/STLFunctionalExtras.h:69:12
#32 0x000056c686e42aba 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) /mlir/lib/Support/ToolUtilities.cpp:30:12
#33 0x000056c678036dc6 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) /mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:784:26
#34 0x000056c6780371c8 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) /mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:830:14
#35 0x000056c678037398 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) /mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:846:10
#36 0x000056c677f058d5 main /mlir/tools/mlir-opt/mlir-opt.cpp:343:33
#37 0x000070a08a42a1ca __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#38 0x000070a08a42a28b call_init ./csu/../csu/libc-start.c:128:20
#39 0x000070a08a42a28b __libc_start_main ./csu/../csu/libc-start.c:347:5
#40 0x000056c677f05745 _start (./mlir-opt+0x481e745)
Segmentation fault
```
#### Reproduction Command
```bash
mlir-opt asyncAPI.mlir -test-vulkan-runner-pipeline
```
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
