| Issue |
177823
|
| Summary |
mlir-opt --nvgpu-optimize-shared-memory crashes on memref with nested vector types
|
| Labels |
new issue
|
| Assignees |
|
| Reporter |
Subway2023
|
# Description
Running mlir-opt with the --nvgpu-optimize-shared-memory pass on a module containing a memref with nested vector elements triggers an assertion failure
# Reproduce
mlir-opt version: 22.1.0-rc1
location: If the **memref.alloc** is commented out, the crash no longer occurs.
```mlir
module {
func.func @main() {
%0 = vector.step : vector<16xindex>
%1 = vector.vscale
scf.execute_region {
scf.yield
}
%alloc = memref.alloc() : memref<16x1xvector<16xf16>, 3>
return
}
}
```
```
mlir-opt --nvgpu-optimize-shared-memory test.mlir
```
```
mlir-opt: /mnt/sdd1/mlir/tool/llvm-project-llvmorg-22.1.0-rc1/mlir/lib/IR/Types.cpp:123: unsigned int mlir::Type::getIntOrFloatBitWidth() const: Assertion `isIntOrFloat() && "only integers and floats have a bitwidth"' 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: /mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt --nvgpu-optimize-shared-memory /mnt/sdd1/sbw/mlir/data/genProgram_1224/250143_94c7/mlir_steps/250143_94c7_5.mlir
#0 0x00005a8e5aecfe32 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0x252fe32)
#1 0x00005a8e5aecc5af llvm::sys::RunSignalHandlers() (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0x252c5af)
#2 0x00005a8e5aecc6fc SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
#3 0x000070e7ee645330 (/lib/x86_64-linux-gnu/libc.so.6+0x45330)
#4 0x000070e7ee69eb2c __pthread_kill_implementation ./nptl/pthread_kill.c:44:76
#5 0x000070e7ee69eb2c __pthread_kill_internal ./nptl/pthread_kill.c:78:10
#6 0x000070e7ee69eb2c pthread_kill ./nptl/pthread_kill.c:89:10
#7 0x000070e7ee64527e raise ./signal/../sysdeps/posix/raise.c:27:6
#8 0x000070e7ee6288ff abort ./stdlib/abort.c:81:7
#9 0x000070e7ee62881b _nl_load_domain ./intl/loadmsgcat.c:1177:9
#10 0x000070e7ee63b517 (/lib/x86_64-linux-gnu/libc.so.6+0x3b517)
#11 0x00005a8e62b1b6b6 mlir::Type::getIntOrFloatBitWidth() const (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0xa17b6b6)
#12 0x00005a8e5c48b533 mlir::nvgpu::optimizeSharedMemoryReadsAndWrites(mlir::Operation*, mlir::Value) (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0x3aeb533)
#13 0x00005a8e5c48cb7a (anonymous namespace)::OptimizeSharedMemoryPass::runOnOperation() OptimizeSharedMemory.cpp:0:0
#14 0x00005a8e6276f756 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0x9dcf756)
#15 0x00005a8e6276fa8e mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0x9dcfa8e)
#16 0x00005a8e627701e2 mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0x9dd01e2)
#17 0x00005a8e62771400 mlir::PassManager::run(mlir::Operation*) (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0x9dd1400)
#18 0x00005a8e5af96767 performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#19 0x00005a8e5af974ea 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&)::'lambda'(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef, llvm::raw_ostream&)>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef const&, llvm::raw_ostream&) MlirOptMain.cpp:0:0
#20 0x00005a8e62b3f731 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) (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0xa19f731)
#21 0x00005a8e5af8e8d0 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (.part.0) MlirOptMain.cpp:0:0
#22 0x00005a8e5af97b70 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0x25f7b70)
#23 0x00005a8e5af97da9 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0x25f7da9)
#24 0x00005a8e5adc9023 main (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0x2429023)
#25 0x000070e7ee62a1ca __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#26 0x000070e7ee62a28b call_init ./csu/../csu/libc-start.c:128:20
#27 0x000070e7ee62a28b __libc_start_main ./csu/../csu/libc-start.c:347:5
#28 0x00005a8e5aeae405 _start (/mnt/sdd1/sbw/mlir/tool/llvmorg-22.1.0-install/bin/mlir-opt+0x250e405)
Aborted (core dumped)
```
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
