| Issue |
185174
|
| Summary |
[MLIR] Crash when lowering gpu function to llvm due to attribute sorting issue
|
| Labels |
mlir
|
| Assignees |
|
| Reporter |
jiayiw2008
|
Hi, I encountered a crash bug. The code example and stack trace are listed below.
Git version: 47766d7f8c397da857d9f78db36c71e8a2ca1fbf
```llvm
module {
gpu.module @CacheSegmentTest {
gpu.func @kernCacheSegmentTest(%data: memref<?xi64>, %dataSize: i32, %trash: memref<?xi32>, %bigDataSize: i32, %hitCount: i32) kernel attributes {linkage = #llvm.linkage<external>} {
%zero_i64 = arith.constant 0 : i64
%_one_i64_ = arith.constant 1 : i64
%zero_index = arith.constant 0 : index
%_one_index_ = arith.constant 1 : index
%shared_counter = gpu.alloc() : memref<4xi16, 3>
%block_xy = gpu.block_id y
%thread_xy = gpu.thread_id y
%block_xy_i32 = arith.index_cast %block_xy : index to i32
%thread_xy_i32 = arith.index_cast %thread_xy : index to i32
%sum_ids = arith.addi %block_xy_i32, %thread_xy_i32 : i32
%mod_sum = arith.remsi %sum_ids, %dataSize : i32
%idx_signed = arith.index_cast %mod_sum : i32 to index
scf.parallel (%i) = (%zero_index) to (%one_index) step (%one_index) {
%loop_id = arith.index_cast %i : index to i32
%addr_i32 = arith.remsi %loop_id, %bigDataSize : i32
%addr = arith.index_cast %addr_i32 : i32 to index
%loaded = memref.load %trash[%addr] : memref<?xi32>
%c2_i32 = arith.constant 2 : i32
%shl = arith.shli %loaded, %c2_i32 : i32
memref.store %shl, %trash[%addr] : memref<?xi32>
}
%zero_idx = arith.constant 0 : index
%_one_idx_ = arith.constant 1 : index
%two_idx = arith.constant 2 : index
%three_idx = arith.constant 3 : index
%simple_constant_0 = arith.constant 0 : i16
%c0 = arith.constant 0 : i64
%ui16_zero_i16 = arith.constant 0 : i16
memref.store %ui16_zero_i16, %shared_counter[%zero_idx] : memref<4xi16, 3>
memref.store %ui16_zero_i16, %shared_counter[%one_idx] : memref<4xi16, 3>
memref.store %ui16_zero_i16, %shared_counter[%two_idx] : memref<4xi16, 3>
memref.store %ui16_zero_i16, %shared_counter[%three_idx] : memref<4xi16, 3>
%hits_tmp = arith.constant 0 : i64
scf.while (%hit_large = %hits_tmp) : (i64) -> i64 {
%hit_i64_trunc = arith.trunci %hit_large : i64 to i32
%cmp = arith.cmpi slt, %hit_i64_trunc, %hitCount : i32
scf.condition(%cmp) %hit_large : i64
} do {
^bb0(%hit_val: i64):
%data_val = memref.load %data[%idx_signed] : memref<?xi64>
%new_data = arith.xori %data_val, %hit_val : i64
memref.store %new_data, %data[%idx_signed] : memref<?xi64>
%four = arith.constant 4 : i32
%counter_slot = arith.remsi %mod_sum, %four : i32
%counter_slot_idx = arith.index_cast %counter_slot : i32 to index
%counter_load = memref.load %shared_counter[%counter_slot_idx] : memref<4xi16, 3>
%counter_slot_i16 = arith.trunci %counter_slot : i32 to i16
%inc_counter = arith.addi %counter_load, %counter_slot_i16 : i16
memref.store %inc_counter, %shared_counter[%counter_slot_idx] : memref<4xi16, 3>
gpu.barrier
%next_hit = arith.addi %hit_val, %one_i64 : i64
scf.yield %next_hit : i64
}
gpu.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 simpleAttributes.mlir -lower-affine -gpu-lower-to-nvvm-pipeline
#0 0x0000575d8b01fbed llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) /llvm/lib/Support/Unix/Signals.inc:880:11
#1 0x0000575d8b02011b PrintStackTraceSignalHandler(void*) /llvm/lib/Support/Unix/Signals.inc:962:1
#2 0x0000575d8b01dfb4 llvm::sys::RunSignalHandlers() /llvm/lib/Support/Signals.cpp:108:5
#3 0x0000575d8b0208b9 SignalHandler(int, siginfo_t*, void*) /llvm/lib/Support/Unix/Signals.inc:448:38
#4 0x00007b061ce45330 (/lib/x86_64-linux-gnu/libc.so.6+0x45330)
#5 0x00007b061ce9eb2c __pthread_kill_implementation ./nptl/pthread_kill.c:44:76
#6 0x00007b061ce9eb2c __pthread_kill_internal ./nptl/pthread_kill.c:78:10
#7 0x00007b061ce9eb2c pthread_kill ./nptl/pthread_kill.c:89:10
#8 0x00007b061ce4527e raise ./signal/../sysdeps/posix/raise.c:27:6
#9 0x00007b061ce288ff abort ./stdlib/abort.c:81:7
#10 0x00007b061ce2881b _nl_load_domain ./intl/loadmsgcat.c:1177:9
#11 0x00007b061ce3b517 (/lib/x86_64-linux-gnu/libc.so.6+0x3b517)
#12 0x0000575d99cc54c3 mlir::DictionaryAttr::sortInPlace(llvm::SmallVectorImpl<mlir::NamedAttribute>&) /mlir/lib/IR/BuiltinAttributes.cpp:129:10
#13 0x0000575d99e35d7a mlir::NamedAttrList::getDictionary(mlir::MLIRContext*) const /mlir/lib/IR/OperationSupport.cpp:56:5
#14 0x0000575d99e26f20 mlir::Operation::create(mlir::OperationState const&) /mlir/lib/IR/Operation.cpp:37:31
#15 0x0000575d99cb7b91 mlir::OpBuilder::create(mlir::OperationState const&) /mlir/lib/IR/Builders.cpp:462:17
#16 0x0000575d985f0f30 mlir::LLVM::LLVMFuncOp::create(mlir::OpBuilder&, mlir::Location, llvm::StringRef, mlir::Type, mlir::LLVM::linkage::Linkage, bool, mlir::LLVM::cconv::CConv, mlir::SymbolRefAttr, llvm::ArrayRef<mlir::NamedAttribute>, llvm::ArrayRef<mlir::DictionaryAttr>, std::optional<unsigned long>) /build/tools/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.cpp.inc:35047:55
#17 0x0000575d8dd65394 mlir::GPUFuncOpLowering::matchAndRewrite(mlir::gpu::GPUFuncOp, mlir::gpu::GPUFuncOpAdaptor, mlir::ConversionPatternRewriter&) const /mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp:234:21
#18 0x0000575d8d942e5b llvm::LogicalResult mlir::ConversionPattern::dispatchTo1To1<mlir::ConvertOpToLLVMPattern<mlir::gpu::GPUFuncOp, false>, mlir::gpu::GPUFuncOp>(mlir::ConvertOpToLLVMPattern<mlir::gpu::GPUFuncOp, false> const&, mlir::gpu::GPUFuncOp, mlir::gpu::GPUFuncOp::GenericAdaptor<llvm::ArrayRef<mlir::ValueRange>>, mlir::ConversionPatternRewriter&) /mlir/include/mlir/Transforms/DialectConversion.h:1069:15
#19 0x0000575d8d942b8e mlir::ConvertOpToLLVMPattern<mlir::gpu::GPUFuncOp, false>::matchAndRewrite(mlir::gpu::GPUFuncOp, mlir::gpu::GPUFuncOpGenericAdaptor<llvm::ArrayRef<mlir::ValueRange>>, mlir::ConversionPatternRewriter&) const /mlir/include/mlir/Conversion/LLVMCommon/Pattern.h:276:12
#20 0x0000575d8d942aa7 mlir::ConvertOpToLLVMPattern<mlir::gpu::GPUFuncOp, false>::matchAndRewrite(mlir::Operation*, llvm::ArrayRef<mlir::ValueRange>, mlir::ConversionPatternRewriter&) const /mlir/include/mlir/Conversion/LLVMCommon/Pattern.h:262:12
#21 0x0000575d997e34a9 mlir::ConversionPattern::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&) const /mlir/lib/Transforms/Utils/DialectConversion.cpp:2410:10
#22 0x0000575d9987b738 mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<llvm::LogicalResult (mlir::Pattern const&)>)::$_0::operator()() const /mlir/lib/Rewrite/PatternApplicator.cpp:223:31
#23 0x0000575d9987b3e5 void llvm::function_ref<void ()>::callback_fn<mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<llvm::LogicalResult (mlir::Pattern const&)>)::$_0>(long) /llvm/include/llvm/ADT/STLFunctionalExtras.h:46:5
#24 0x0000575d8b045fb9 llvm::function_ref<void ()>::operator()() const /llvm/include/llvm/ADT/STLFunctionalExtras.h:69:5
#25 0x0000575d9987cf3b void mlir::MLIRContext::executeAction<mlir::ApplyPatternAction, mlir::Pattern const&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pattern const&) /mlir/include/mlir/IR/MLIRContext.h:291:3
#26 0x0000575d998797ca mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<llvm::LogicalResult (mlir::Pattern const&)>) /mlir/lib/Rewrite/PatternApplicator.cpp:242:9
#27 0x0000575d997f0088 (anonymous namespace)::OperationLegalizer::legalizeWithPattern(mlir::Operation*) /mlir/lib/Transforms/Utils/DialectConversion.cpp:2831:21
#28 0x0000575d997e41a2 (anonymous namespace)::OperationLegalizer::legalize(mlir::Operation*) /mlir/lib/Transforms/Utils/DialectConversion.cpp:2609:17
#29 0x0000575d997e3a72 mlir::OperationConverter::convert(mlir::Operation*, bool) /mlir/lib/Transforms/Utils/DialectConversion.cpp:3309:26
#30 0x0000575d997e4b3d llvm::LogicalResult mlir::OperationConverter::legalizeOperations<mlir::OperationConverter::applyConversion(llvm::ArrayRef<mlir::Operation*>)::$_0>(llvm::ArrayRef<mlir::Operation*>, mlir::OperationConverter::applyConversion(llvm::ArrayRef<mlir::Operation*>)::$_0, bool) /mlir/lib/Transforms/Utils/DialectConversion.cpp:3410:16
#31 0x0000575d997e46ee mlir::OperationConverter::applyConversion(llvm::ArrayRef<mlir::Operation*>) /mlir/lib/Transforms/Utils/DialectConversion.cpp:3456:26
#32 0x0000575d997f556c applyConversion(llvm::ArrayRef<mlir::Operation*>, mlir::ConversionTarget const&, mlir::FrozenRewritePatternSet const&, mlir::ConversionConfig, (anonymous namespace)::OpConversionMode)::$_0::operator()() const /mlir/lib/Transforms/Utils/DialectConversion.cpp:4179:30
#33 0x0000575d997f54d5 void llvm::function_ref<void ()>::callback_fn<applyConversion(llvm::ArrayRef<mlir::Operation*>, mlir::ConversionTarget const&, mlir::FrozenRewritePatternSet const&, mlir::ConversionConfig, (anonymous namespace)::OpConversionMode)::$_0>(long) /llvm/include/llvm/ADT/STLFunctionalExtras.h:46:5
#34 0x0000575d8b045fb9 llvm::function_ref<void ()>::operator()() const /llvm/include/llvm/ADT/STLFunctionalExtras.h:69:5
#35 0x0000575d998118d3 void mlir::MLIRContext::executeAction<ApplyConversionAction>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>) /mlir/include/mlir/IR/MLIRContext.h:291:3
#36 0x0000575d997e82e8 applyConversion(llvm::ArrayRef<mlir::Operation*>, mlir::ConversionTarget const&, mlir::FrozenRewritePatternSet const&, mlir::ConversionConfig, (anonymous namespace)::OpConversionMode) /mlir/lib/Transforms/Utils/DialectConversion.cpp:4183:1
#37 0x0000575d997e819c mlir::applyPartialConversion(llvm::ArrayRef<mlir::Operation*>, mlir::ConversionTarget const&, mlir::FrozenRewritePatternSet const&, mlir::ConversionConfig) /mlir/lib/Transforms/Utils/DialectConversion.cpp:4192:10
#38 0x0000575d997e8385 mlir::applyPartialConversion(mlir::Operation*, mlir::ConversionTarget const&, mlir::FrozenRewritePatternSet const&, mlir::ConversionConfig) /mlir/lib/Transforms/Utils/DialectConversion.cpp:4199:10
#39 0x0000575d8d946078 (anonymous namespace)::LowerGpuOpsToNVVMOpsPass::runOnOperation() /mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp:428:13
#40 0x0000575d9993d1d4 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_3::operator()() const /mlir/lib/Pass/Pass.cpp:0:19
#41 0x0000575d9993d175 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
#42 0x0000575d8b045fb9 llvm::function_ref<void ()>::operator()() const /llvm/include/llvm/ADT/STLFunctionalExtras.h:69:5
#43 0x0000575d9994067b 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
#44 0x0000575d9993461e mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /mlir/lib/Pass/Pass.cpp:619:17
#45 0x0000575d99934d3a 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
#46 0x0000575d9993efe8 mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_12::operator()(mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo&) const /mlir/lib/Pass/Pass.cpp:1001:36
#47 0x0000575d9993ee80 auto void mlir::parallelForEach<__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_12>(mlir::MLIRContext*, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_12&&)::'lambda'(__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>&&)::operator()<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo&>(__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>&&) const /mlir/include/mlir/IR/Threading.h:120:56
#48 0x0000575d9993ec68 llvm::LogicalResult mlir::failableParallelForEach<__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>, void mlir::parallelForEach<__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo>>>, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_12>(mlir::MLIRContext*, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::Op<truncated>Please see the issue for the entire body.
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
