| Issue |
182745
|
| Summary |
[MLIR] Segmentation fault in LaunchOp::verifyRegions in a gpu program
|
| Labels |
mlir
|
| Assignees |
|
| Reporter |
jiayiw2008
|
Hi, I've found a segmentation fault in the following code example.
```llvm
module {
func.func @host_launch_hgemm_sliced_k_f16_kernel(%arg0: memref<?x?xf16>, %arg1: memref<?x?xf16>, %arg2: memref<?x?xf16>) {
%c0 = arith.constant 0 : index
%c1 = arith.constant 1 : index
%c2 = arith.constant 2 : index
%M = memref.dim %arg0, %c0 : memref<?x?xf16>
%N = memref.dim %arg1, %c1 : memref<?x?xf16>
%grid_x = arith.ceildivui %N, %c2 : index
%grid_y = arith.ceildivui %M, %c2 : index
gpu.launch blocks(%bx, %by, %bz) in (%grid_sz_x = %grid_x, %grid_sz_y = %grid_y, %grid_sz_z = %c1)
threads(%tx, %ty, %tz) in (%block_sz_x = %c2, %block_sz_y = %c2, %block_sz_z = %c1)
workgroup(%s_a : memref<2x2xf16, 3>, %s_b : memref<2x2xf16, 3>) {
%K = memref.dim %arg0, %c1 : memref<?x?xf16>
%c2_1 = arith.constant 2 : index
%tid = arith.muli %ty, %tx : index
%load_smem_a_m = arith.remui %tid, %c2_1 : index
%load_smem_a_k = arith.divui %tid, %c2_1 : index
%load_smem_b_k = arith.remui %tid, %c2_1 : index
%load_smem_b_n = arith.divui %tid, %c2_1 : index
%load_gmem_a_m = arith.addi %by, %load_smem_a_m : index
%load_gmem_b_n = arith.addi %bx, %load_smem_b_n : index
%cond_a = arith.cmpi ult, %load_gmem_a_m, %M : index
%cond_b = arith.cmpi ult, %load_gmem_b_n, %N : index
%cond = arith.andi %cond_a, %cond_b : i1
scf.if %cond {
%c0_2 = arith.constant 0 : index
%c2_3 = arith.constant 2 : index
%num_bk = arith.ceildivui %K, %c2_3 : index
%sum = arith.constant 0.000000e+00 : f16
%outer = scf.for %bk = %c0_2 to %num_bk step %c1 iter_args(%acc = %sum) -> (f16) {
%load_gmem_a_k = arith.addi %bk, %load_smem_a_k : index
%a_val = memref.load %arg0[%load_gmem_a_m, %load_gmem_a_k] : memref<?x?xf16>
memref.store %a_val, %s_a[%load_smem_a_m, %load_smem_a_k] : memref<2x2xf16, 3>
%load_gmem_b_k = arith.addi %bk, %load_smem_b_k : index
%b_val = memref.load %arg1[%load_gmem_b_k, %load_gmem_b_n] : memref<?x?xf16>
memref.store %b_val, %s_b[%load_smem_b_k, %load_smem_b_n] : memref<2x2xf16, 3>
gpu.barrier
%c0_4 = arith.constant 0 : index
%inner = scf.for %k = %c0_4 to %c2_3 step %c1 iter_args(%inner_acc = %acc) -> (f16) {
%a_smem = memref.load %s_a[%load_smem_a_m, %k] : memref<2x2xf16, 3>
%b_smem = memref.load %s_b[%k, %load_smem_b_n] : memref<2x2xf16, 3>
%prod = arith.mulf %a_smem, %b_smem : f16
%new_acc = arith.addf %inner_acc, %prod : f16
scf.yield %new_acc : f16
}
scf.yield %inner : f16
}
memref.store %outer, %arg2[%load_gmem_a_m, %load_gmem_b_n] : memref<?x?xf16>
}
gpu.terminator
}
return
}
}
```
#### 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 --test-ir-visitors --gpu-kernel-outlining --convert-gpu-to-nvvm
#0 0x00006265c4deb1fd llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) /llvm/lib/Support/Unix/Signals.inc:880:11
#1 0x00006265c4deb72b PrintStackTraceSignalHandler(void*) /llvm/lib/Support/Unix/Signals.inc:962:1
#2 0x00006265c4de95c4 llvm::sys::RunSignalHandlers() /llvm/lib/Support/Signals.cpp:108:5
#3 0x00006265c4debec9 SignalHandler(int, siginfo_t*, void*) /llvm/lib/Support/Unix/Signals.inc:448:38
#4 0x00007ec8e8245330 (/lib/x86_64-linux-gnu/libc.so.6+0x45330)
#5 0x00006265d1b4952e verifyAttributions(mlir::Operation*, llvm::ArrayRef<mlir::BlockArgument>, mlir::gpu::AddressSpace) /mlir/lib/Dialect/GPU/IR/GPUDialect.cpp:567:16
#6 0x00006265d1b491b1 mlir::gpu::LaunchOp::verifyRegions() /mlir/lib/Dialect/GPU/IR/GPUDialect.cpp:902:14
#7 0x00006265d1ce58d3 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*) /mlir/include/mlir/IR/OpDefinition.h:2065:39
#8 0x00006265c5170265 llvm::LogicalResult llvm::detail::UniqueFunctionBase<llvm::LogicalResult, mlir::Operation*>::CallImpl<llvm::LogicalResult (* const)(mlir::Operation*)>(void*, mlir::Operation*) /llvm/include/llvm/ADT/FunctionExtras.h:212:12
#9 0x00006265c5170037 llvm::unique_function<llvm::LogicalResult (mlir::Operation*) const>::operator()(mlir::Operation*) const /llvm/include/llvm/ADT/FunctionExtras.h:390:12
#10 0x00006265d1ce4066 mlir::RegisteredOperationName::Model<mlir::gpu::LaunchOp>::verifyRegionInvariants(mlir::Operation*) /mlir/include/mlir/IR/OperationSupport.h:558:14
#11 0x00006265d399efa6 mlir::OperationName::verifyRegionInvariants(mlir::Operation*) const /mlir/include/mlir/IR/OperationSupport.h:317:23
#12 0x00006265d399c2df (anonymous namespace)::OperationVerifier::verifyOnExit(mlir::Operation&) /mlir/lib/IR/Verifier.cpp:236:48
#13 0x00006265d399c020 _ZZN12_GLOBAL__N_117OperationVerifier15verifyOperationERN4mlir9OperationEENK3$_0clIS2_EEDaPT_ /mlir/lib/IR/Verifier.cpp:287:53
#14 0x00006265d399bf55 _ZZN12_GLOBAL__N_117OperationVerifier15verifyOperationERN4mlir9OperationEENK3$_1clIZNS0_15verifyOperationES3_E3$_0EEDaOT_N4llvm12PointerUnionIJPS2_PNS1_5BlockEEEE /mlir/lib/IR/Verifier.cpp:276:16
#15 0x00006265d399b81a (anonymous namespace)::OperationVerifier::verifyOperation(mlir::Operation&) /mlir/lib/IR/Verifier.cpp:287:15
#16 0x00006265d399b681 (anonymous namespace)::OperationVerifier::verifyOpAndDominance(mlir::Operation&) /mlir/lib/IR/Verifier.cpp:81:14
#17 0x00006265d399b632 mlir::verify(mlir::Operation*, bool) /mlir/lib/IR/Verifier.cpp:425:19
#18 0x00006265d3789028 verifyOpAndAdjustFlags(mlir::Operation*, mlir::OpPrintingFlags) /mlir/lib/IR/AsmPrinter.cpp:2099:14
#19 0x00006265d3788f57 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*) /mlir/lib/IR/AsmPrinter.cpp:2111:15
#20 0x00006265d379014b mlir::Block::printAsOperand(llvm::raw_ostream&, bool) /mlir/lib/IR/AsmPrinter.cpp:4206:12
#21 0x00006265c5c68e4b printBlock(mlir::Block*) /mlir/test/lib/IR/TestVisitors.cpp:24:3
#22 0x00006265c5c6aeb2 testNoSkipErasureCallbacks(mlir::Operation*)::$_1::operator()(mlir::Block*) const /mlir/test/lib/IR/TestVisitors.cpp:186:7
#23 0x00006265c5c6ae5d void llvm::function_ref<void (mlir::Block*)>::callback_fn<testNoSkipErasureCallbacks(mlir::Operation*)::$_1>(long, mlir::Block*) /llvm/include/llvm/ADT/STLFunctionalExtras.h:46:5
#24 0x00006265c5011551 llvm::function_ref<void (mlir::Block*)>::operator()(mlir::Block*) const /llvm/include/llvm/ADT/STLFunctionalExtras.h:69:5
#25 0x00006265c5011398 void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Block*)>, mlir::WalkOrder) /mlir/include/mlir/IR/Visitors.h:123:5
#26 0x00006265c5011374 void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Block*)>, mlir::WalkOrder) /mlir/include/mlir/IR/Visitors.h:127:27
#27 0x00006265c5c6ae02 std::enable_if<llvm::is_one_of<mlir::Block*, mlir::Operation*, mlir::Region*, mlir::Block*>::value, void>::type mlir::detail::walk<(mlir::WalkOrder)1, mlir::ForwardIterator, testNoSkipErasureCallbacks(mlir::Operation*)::$_1&, mlir::Block*, void>(mlir::Operation*, testNoSkipErasureCallbacks(mlir::Operation*)::$_1&) /mlir/include/mlir/IR/Visitors.h:278:3
#28 0x00006265c5c6ac9d std::enable_if<llvm::function_traits<std::decay<testNoSkipErasureCallbacks(mlir::Operation*)::$_1&>::type>::num_args == 1, void>::type mlir::Operation::walk<(mlir::WalkOrder)1, mlir::ForwardIterator, testNoSkipErasureCallbacks(mlir::Operation*)::$_1&, void>(testNoSkipErasureCallbacks(mlir::Operation*)::$_1&) /mlir/include/mlir/IR/Operation.h:798:5
#29 0x00006265c5c689be testNoSkipErasureCallbacks(mlir::Operation*) /mlir/test/lib/IR/TestVisitors.cpp:222:3
#30 0x00006265c5c6832d (anonymous namespace)::TestIRVisitorsPass::runOnOperation() /mlir/test/lib/IR/TestVisitors.cpp:293:3
#31 0x00006265d3465734 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_3::operator()() const /mlir/lib/Pass/Pass.cpp:0:19
#32 0x00006265d34656d5 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
#33 0x00006265c4e115c9 llvm::function_ref<void ()>::operator()() const /llvm/include/llvm/ADT/STLFunctionalExtras.h:69:5
#34 0x00006265d3468bdb 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
#35 0x00006265d345cb7e mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /mlir/lib/Pass/Pass.cpp:619:17
#36 0x00006265d345d29a 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
#37 0x00006265d3462b7d mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) /mlir/lib/Pass/Pass.cpp:1123:10
#38 0x00006265d3462106 mlir::PassManager::run(mlir::Operation*) /mlir/lib/Pass/Pass.cpp:1097:60
#39 0x00006265c4ef5b4b performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) /mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:591:17
#40 0x00006265c4ef5221 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:673:12
#41 0x00006265c4ef4edc 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:771:12
#42 0x00006265c4ef4de6 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
#43 0x00006265d39c788a 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
#44 0x00006265d39c6c4a 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
#45 0x00006265c4ef0b66 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:776:26
#46 0x00006265c4ef0f68 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) /mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:822:14
#47 0x00006265c4ef1138 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) /mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:838:10
#48 0x00006265c4dbfde5 main /mlir/tools/mlir-opt/mlir-opt.cpp:347:33
#49 0x00007ec8e822a1ca __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#50 0x00007ec8e822a28b call_init ./csu/../csu/libc-start.c:128:20
#51 0x00007ec8e822a28b __libc_start_main ./csu/../csu/libc-start.c:347:5
#52 0x00006265c4dbfc55 _start (/mnt/raid/build_all/bin/mlir-opt+0x4742c55)
Segmentation fault
```
Version: 7dbfc2bd0078b36e4f9ed14dd25c4d8e97b3ae77
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
