Open
Description
Test on commit: 6548b63
steps to reproduce:
mlir-opt test.mlir --gpu-module-to-binary=format=%gpu_compilation_format
test case 1:
module {
func.func @transpose_nofold_non_cst_input(%arg0: tensor<2x3xf32> {bufferization.access = "read"}) -> tensor<3x2xf32> {
%cst = arith.constant dense<[1, 0]> : tensor<2xi32>
%0 = tosa.transpose %arg0, %cst {__inplace_operands_attr__ = ["false", "false"]} : (tensor<2x3xf32>, tensor<2xi32>) -> tensor<3x2xf32>
return {__inplace_operands_attr__ = ["true"]} %0 : tensor<3x2xf32>
}
}
test case 2:
module {
func.func @test_maxpool2d_stride_y(%arg0: tensor<1x32x32x8xf32>) -> tensor<1x32x32x8xf32> {
%0 = "tosa.max_pool2d"(%arg0) {kernel = array<i64: 1, 1>, pad = array<i64: 4, 4, 4, 4>, stride = array<i64: 8193, 1>} : (tensor<1x32x32x8xf32>) -> tensor<1x32x32x8xf32>
return %0 : tensor<1x32x32x8xf32>
}
}
crash trace:
/usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/optional:439: _Tp &std::_Optional_base_impl<mlir::gpu::CompilationTarget, std::_Optional_base<mlir::gpu::CompilationTarget>>::_M_get() [_Tp = mlir::gpu::CompilationTarget, _Dp = std::_Optional_base<mlir::gpu::CompilationTarget>]: Assertion 'this->_M_is_engaged()' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0. Program arguments: ./mlir-opt /home/workdir/test.mlir --gpu-module-to-binary=format=%gpu_compilation_format
#0 0x00005da9823233b8 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (./mlir-opt+0x10723b8)
#1 0x00005da982320ede llvm::sys::RunSignalHandlers() (./mlir-opt+0x106fede)
#2 0x00005da982323dc8 SignalHandler(int) Signals.cpp:0:0
#3 0x00007e2f641f5520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
#4 0x00007e2f642499fc pthread_kill (/lib/x86_64-linux-gnu/libc.so.6+0x969fc)
#5 0x00007e2f641f5476 gsignal (/lib/x86_64-linux-gnu/libc.so.6+0x42476)
#6 0x00007e2f641db7f3 abort (/lib/x86_64-linux-gnu/libc.so.6+0x287f3)
#7 0x00005da98230fcaf (./mlir-opt+0x105ecaf)
#8 0x00005da982af29a8 (anonymous namespace)::GpuModuleToBinaryPass::runOnOperation() ModuleToBinary.cpp:0:0
#9 0x00005da98559f6df mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (./mlir-opt+0x42ee6df)
#10 0x00005da98559fec2 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (./mlir-opt+0x42eeec2)
#11 0x00005da9855a266e mlir::PassManager::run(mlir::Operation*) (./mlir-opt+0x42f166e)
#12 0x00005da98559b192 performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#13 0x00005da98559adfb llvm::LogicalResult llvm::function_ref<llvm::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, 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::raw_ostream&) MlirOptMain.cpp:0:0
#14 0x00005da985644ae5 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::raw_ostream&)>, llvm::raw_ostream&, llvm::StringRef, llvm::StringRef) (./mlir-opt+0x4393ae5)
#15 0x00005da985595ff2 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (./mlir-opt+0x42e4ff2)
#16 0x00005da9855962a3 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (./mlir-opt+0x42e52a3)
#17 0x00005da9855964b2 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (./mlir-opt+0x42e54b2)
#18 0x00005da9823025d7 main (./mlir-opt+0x10515d7)
#19 0x00007e2f641dcd90 (/lib/x86_64-linux-gnu/libc.so.6+0x29d90)
#20 0x00007e2f641dce40 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x29e40)
#21 0x00005da982302145 _start (./mlir-opt+0x1051145)
Aborted (core dumped)