llvm / llvm/llvm-project

[MLIR] Crash in `deriveStaticUpperBound` when lowering nested `scf.parallel` loops

Open
#224,211 1 comment 0 reactions 0 assignees View on GitHub
crash mlir
Dominant language
LLVM
Stars
40.5k
Forks
18.7k
PR merge metrics
PR metrics pending

Description

### Description

I encountered a crash when lowering a `nested scf.parallel` loop to `GPU`. The crash occurs in `deriveStaticUpperBound()` in `SCFToGPU.cpp`. The assertion indicates that `dyn_cast` is called on a non-existent mlir::Value.

### Tested commit

[3b8ee89](https://github.com/llvm/llvm-project/commit/3b8ee89006f5a228fce75528f8dd80b4229abcc1)

### Test case

```mlir
module attributes {llvm.data_layout = ""} {
emitc.include <"cstdlib">
func.func @main() {
%2 = "emitc.constant"() <{value = 64 : index}> : () -> !emitc.size_t
%3 = "emitc.constant"() <{value = 1024 : index}> : () -> index
%4 = llvm.mlir.constant(1 : i64) : i64
%5 = llvm.mlir.constant(0 : i64) : i64
%7 = llvm.mlir.constant(16 : i64) : i64
%10 = builtin.unrealized_conversion_cast %7 : i64 to index
%11 = builtin.unrealized_conversion_cast %4 : i64 to index
%12 = builtin.unrealized_conversion_cast %5 : i64 to index
%c0 = arith.constant 0 : index
%c1 = arith.constant 1 : index
%19 = arith.muli %11, %c1 : index
scf.parallel (%arg2, %arg3) = (%12, %12) to (%10, %10) step (%19, %19) {
%c-1 = arith.constant -1 : index
%25 = arith.minsi %19, %c-1 : index
scf.parallel (%arg5) = (%c0) to (%25) step (%11) {
scf.reduce
}
scf.reduce
}
return
}
}
```

### Steps to reproduce

```
mlir-opt test.mlir -gpu-map-parallel-loops -convert-parallel-loops-to-gpu
```

### stack trace

```
mlir-opt: /home/workdir/llvm-project-latest/llvm/include/llvm/Support/Casting.h:644: decltype(auto) llvm::dyn_cast(const From &) [To = mlir::OpResult, From = mlir::Value]: Assertion `detail::isPresent(Val) && "dyn_cast on a non-existent value"' failed.
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 test.mlir -gpu-map-parallel-loops -convert-parallel-loops-to-gpu
#0 0x00005602e431fd0b llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (../llvm-project-latest/build/bin/mlir-opt+0x1fc2d0b)
#1 0x00005602e431cb11 llvm::sys::RunSignalHandlers() (../llvm-project-latest/build/bin/mlir-opt+0x1fbfb11)
#2 0x00005602e4320a3b SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
#3 0x00007f1646305520 (/usr/lib/x86_64-linux-gnu/libc.so.6+0x42520)
#4 0x00007f16463599fc pthread_kill (/usr/lib/x86_64-linux-gnu/libc.so.6+0x969fc)
#5 0x00007f1646305476 gsignal (/usr/lib/x86_64-linux-gnu/libc.so.6+0x42476)
#6 0x00007f16462eb7f3 abort (/usr/lib/x86_64-linux-gnu/libc.so.6+0x287f3)
#7 0x00007f16462eb71b (/usr/lib/x86_64-linux-gnu/libc.so.6+0x2871b)
#8 0x00007f16462fce96 (/usr/lib/x86_64-linux-gnu/libc.so.6+0x39e96)
#9 0x00005602ec863e91 (../llvm-project-latest/build/bin/mlir-opt+0xa506e91)
#10 0x00005602e5ee486b deriveStaticUpperBound(mlir::Value, mlir::PatternRewriter&) SCFToGPU.cpp:0:0
#11 0x00005602e5ee47f3 deriveStaticUpperBound(mlir::Value, mlir::PatternRewriter&) SCFToGPU.cpp:0:0
#12 0x00005602e5ee33e2 processParallelLoop(mlir::scf::ParallelOp, mlir::gpu::LaunchOp, mlir::IRMapping&, llvm::SmallVectorImpl&, llvm::DenseMap, llvm::detail::DenseMapPair>&, mlir::PatternRewriter&) SCFToGPU.cpp:0:0
#13 0x00005602e5ee20a3 (anonymous namespace)::ParallelToGpuLaunchLowering::matchAndRewrite(mlir::scf::ParallelOp, mlir::PatternRewriter&) const SCFToGPU.cpp:0:0
#14 0x00005602ec4c46c8 void llvm::function_ref::callback_fn, llvm::function_ref, llvm::function_ref)::$_0>(long) PatternApplicator.cpp:0:0
#15 0x00005602ec4c0e2a mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref, llvm::function_ref, llvm::function_ref) (../llvm-project-latest/build/bin/mlir-opt+0xa163e2a)
#16 0x00005602ec476f62 (anonymous namespace)::OperationLegalizer::legalize(mlir::Operation*) DialectConversion.cpp:0:0
#17 0x00005602ec476750 mlir::OperationConverter::convert(mlir::Operation*, bool) (../llvm-project-latest/build/bin/mlir-opt+0xa119750)
#18 0x00005602ec477c5e mlir::OperationConverter::applyConversion(llvm::ArrayRef) (../llvm-project-latest/build/bin/mlir-opt+0xa11ac5e)
#19 0x00005602ec4864c9 void llvm::function_ref::callback_fn, mlir::ConversionTarget const&, mlir::FrozenRewritePatternSet const&, mlir::ConversionConfig, (anonymous namespace)::OpConversionMode)::$_0>(long) DialectConversion.cpp:0:0
#20 0x00005602ec47d927 applyConversion(llvm::ArrayRef, mlir::ConversionTarget const&, mlir::FrozenRewritePatternSet const&, mlir::ConversionConfig, (anonymous namespace)::OpConversionMode) DialectConversion.cpp:0:0
#21 0x00005602ec47da0b mlir::applyPartialConversion(mlir::Operation*, mlir::ConversionTarget const&, mlir::FrozenRewritePatternSet const&, mlir::ConversionConfig) (../llvm-project-latest/build/bin/mlir-opt+0xa120a0b)
#22 0x00005602e5edf962 (anonymous namespace)::ParallelLoopToGpuPass::runOnOperation() SCFToGPUPass.cpp:0:0
#23 0x00005602ec51f3ee mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (../llvm-project-latest/build/bin/mlir-opt+0xa1c23ee)
#24 0x00005602ec520364 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (../llvm-project-latest/build/bin/mlir-opt+0xa1c3364)
#25 0x00005602ec5270fa mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) (../llvm-project-latest/build/bin/mlir-opt+0xa1ca0fa)
#26 0x00005602ec526a5d mlir::PassManager::run(mlir::Operation*) (../llvm-project-latest/build/bin/mlir-opt+0xa1c9a5d)
#27 0x00005602e43cb58f performActions(llvm::raw_ostream&, std::shared_ptr const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#28 0x00005602e43ca761 llvm::LogicalResult llvm::function_ref>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>::callback_fn>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&)::$_0>(long, std::unique_ptr>, llvm::MemoryBufferRef const&, llvm::raw_ostream&) MlirOptMain.cpp:0:0
#29 0x00005602ec879ec6 mlir::splitAndProcessBuffer(std::unique_ptr>, llvm::function_ref>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>, llvm::raw_ostream&, llvm::StringRef, llvm::StringRef) (../llvm-project-latest/build/bin/mlir-opt+0xa51cec6)
#30 0x00005602e43c0ba6 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (../llvm-project-latest/build/bin/mlir-opt+0x2063ba6)
#31 0x00005602e43c0f7a mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (../llvm-project-latest/build/bin/mlir-opt+0x2063f7a)
#32 0x00005602e43c11eb mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (../llvm-project-latest/build/bin/mlir-opt+0x20641eb)
#33 0x00005602e43057cd main (../llvm-project-latest/build/bin/mlir-opt+0x1fa87cd)
#34 0x00007f16462ecd90 (/usr/lib/x86_64-linux-gnu/libc.so.6+0x29d90)
#35 0x00007f16462ece40 __libc_start_main (/usr/lib/x86_64-linux-gnu/libc.so.6+0x29e40)
#36 0x00005602e4305325 _start (../llvm-project-latest/build/bin/mlir-opt+0x1fa8325)
Aborted (core dumped)
```

Contributor guide

Open the contributing guide

Research direction

Start with deriveStaticUpperBound and processParallelLoop in SCFToGPU.cpp, then reproduce the crash using the supplied MLIR input and the mlir-opt command. Trace the nested scf.parallel lowering path and use the provided case as regression coverage; done means the command no longer asserts during GPU lowering.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
compilers
Issue type
Bug
Difficulty
3/5
Estimated time
1-2 days
Activity status
Active
Clarity
Clearly specified
Newbie friendliness
66/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.