<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=https://github.com/llvm/llvm-project/issues/116346>116346</a>
        </td>
    </tr>

    <tr>
        <th>Summary</th>
        <td>
            [MLIR]-gpu-module-to-binary="format=%gpu_compilation_format" triggers Assertion Failure 'this->_M_is_engaged()'
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            mlir
      </td>
    </tr>

    <tr>
      <th>Assignees</th>
      <td>
      </td>
    </tr>

    <tr>
      <th>Reporter</th>
      <td>
          xisang0
      </td>
    </tr>
</table>

<pre>
    Test on commit: https://github.com/llvm/llvm-project/commit/6548b6354d1d990e1c98736f5e7c3de876bedc8e
steps to reproduce:
```
mlir-opt test.mlir --gpu-module-to-binary="format=%gpu_compilation_format"
```
test case:
```
module {
  func.func @main(%arg0: tensor<13x21x3xf32>) -> () {
    %ceil_result = tosa.ceil %arg0 : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
    return
  }
}
```
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 0x0000591d42b6e3b8 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (./mlir-opt+0x10723b8)
 #1 0x0000591d42b6bede llvm::sys::RunSignalHandlers() (./mlir-opt+0x106fede)
 #2 0x0000591d42b6edc8 SignalHandler(int) Signals.cpp:0:0
 #3 0x00007e0f5788d520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
 #4 0x00007e0f578e19fc pthread_kill (/lib/x86_64-linux-gnu/libc.so.6+0x969fc)
 #5 0x00007e0f5788d476 gsignal (/lib/x86_64-linux-gnu/libc.so.6+0x42476)
 #6 0x00007e0f578737f3 abort (/lib/x86_64-linux-gnu/libc.so.6+0x287f3)
 #7 0x0000591d42b5acaf (./mlir-opt+0x105ecaf)
 #8 0x0000591d4333d9a8 (anonymous namespace)::GpuModuleToBinaryPass::runOnOperation() ModuleToBinary.cpp:0:0
 #9 0x0000591d45dea6df mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (./mlir-opt+0x42ee6df)
#10 0x0000591d45deaec2 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (./mlir-opt+0x42eeec2)
#11 0x0000591d45ded66e mlir::PassManager::run(mlir::Operation*) (./mlir-opt+0x42f166e)
#12 0x0000591d45de6192 performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#13 0x0000591d45de5dfb 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 0x0000591d45e8fae5 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 0x0000591d45de0ff2 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (./mlir-opt+0x42e4ff2)
#16 0x0000591d45de12a3 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (./mlir-opt+0x42e52a3)
#17 0x0000591d45de14b2 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (./mlir-opt+0x42e54b2)
#18 0x0000591d42b4d5d7 main (./mlir-opt+0x10515d7)
#19 0x00007e0f57874d90 (/lib/x86_64-linux-gnu/libc.so.6+0x29d90)
#20 0x00007e0f57874e40 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x29e40)
#21 0x0000591d42b4d145 _start (./mlir-opt+0x1051145)
Aborted (core dumped)
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzcWF1z4yoS_TX4hbJLAn0--MGJx7tTNamkZvKuQqIls4NAC-je5N9vgRzbcuJJcu9-1N6UY0sCTp_uPg0IZq3oFMAapTco3S7Y6PbarJ-EZaqLFrXmz-tHsA5rhRvd98IhusF75waL6AaRHSK7Trj9WK8a3SOyk_K3l5_lYPQ_oHGI7A5DyS5Lk6LOaJrwmJdlBHFTFjnN2hTyhnIo8qwG3hSAoi2KNtbBYLHT2MBgNB8b8EZDE8qiwyfc9lKYpR4cdmDdyt_h5bIbxmWv-Shh6fSyFoqZZ0S3iJBWm565cJ12w1g1uh-EZE5oVR3aCHnTkMfHDbPXmQSDGOU30z3G7aialf_CKIl6JhQiBSIpM13kg-lAWW0QvY3pE4mf6FNLCaJfECnxEtEvOPQuzwExRiRtQMjKgB2lw4husdOWrfxDfMDGHhyR4j38a-1HYwbcaNTLPcq3B7-PF_MANIbZPXaGXc8WIrvRGi8TUXsBNQ0iu6ciq7JkKYUan5adGhHZxTEiu9Xq9ZdQjRw5eGUhchM-U2c9-BwyiegmoaWPQPU4YEQy67hnQzfV_aFLVTMLlegHieitV8zU3g3jdHF70sQjMx04RG7xFZiPItAv4RMA7qqAOSU3vQk86RZ_lEm1nfr_eUbp1sdpYy0Y34ARyd1eWC-O6q4StgLVsQ74xBWRHLdMSOCrKZkP375sfnzBdqx74TDD9dj5etXG-cr99FQhrB3BIrLDTHF8SDR2e8CTsGrW_AziOpj_4VjzE_OxH45qi1Y4_D0Y3RnWY2a6sQflPA_s5XOcLRDZ7XXvdfS7Nj-58Jr8wAzygeljqhZEaISjpyiKorSMeULqDGhd4OB6yIh9ttPFgxHKBWceQ-mQ4tTJsN8rbZ0B1iOS-ewL5YJuSHHuDyI30VMc5YTWIVVHEvEFiRo4vEXi-6h-iE4x-XemuARjX_T5pp2sBV-DZ3bIpbO8KfAMEZHiwH16bFfN4DMXhf8jED0A5RC1aV4UPCXRNBMe5ow3Zgsp6mZl9SoL9BKSkmhGLpljQly2DR7c3gDj1U8h5ecMlFnZNjMD6SXpJM9wZ4OfnyWf5NkMO5tj5zRvKWa1r7JPIZMib-kMOZ_nLGUNa68kPIWGtbPBxflgSikvWeEHM6XVc69HixXrwQ5B0eUksr8N412oqUd9EyrqgdmD_syo7tX9ACbU0kF8885v66U855FyYBlvz6ZSDo4JOV3fD4_6fvA2N5wNTpujaUSK05BAimx8sZ0enlG7aNkoJp-tsHdMsc7L_BbXWkv_O6qwx-K_KtqEAGT8FFtfs9GlT9CQT_v0IAaQQsHMt6nrkWr2b3fyfJA39VVZZ8IcrM0r1HmHyfLUcOD-wAwo91W1GjdaWRcQrscRGjKLY3wRR55lcGH-xaE3lTCPxzW7bZxlMLNLLuxmcUnwAMavEJvGA9p3Zvnj2m73zACvBuf3aacRP_RoGrjrjN_HHUJzkcy7b1-_32rl4Mm9CvydFOZ-cHdMqFutWtGdYZT4rPV1zXn_6IV_KW_rs0Xlm-5Ew-T3aYt6eu53wmGlNNDOvJkPQKQ4uj8q8c8RXrl_B702zzdj205aPPbn0LJRuoqDBAfXx0zbMXKLryehPG7YGial33xUrZptrs7i9NF0_sf9ObHbCiahcd-hE9aZ59f6-JUGpi6IJFUUgAupVfff9eQXmXlfoclMoVC0DNIz3-0ghdso_mB0A9a-0P2f6O7_sCp-2eO87YczQnXfob36-MqkSkvKIJ1NqunFpBO1LXlbzn-1Yry23iVtO1_vsosQxYTRqyE6LNfNnk0L8-azqXvPv6u8U8LojHd-yTupr6f2M7z_KMGknge2mO-UE57yHPdMqGvb5Tjl-QyhvNjFJ7z85GsNKXkZnWOS6BITkghXlR9WWceMq14ofsYKJHMr8aXvcZLiCf-q93Fyqt2Nf1UB7vs22kB4YQd-MvFyQLTga8pLWrIFrOOcxlkW0YQu9us0SilrUspoSeOoAABGeUrjOioSKEm9EGsSkSSO4zQqkjShqwxS2rKojdMiKxnjKImgZ0KuvExW2nSLcNawjuOMJtlCshqkDQehhATJEILS7cKswwlFPXYWJZEU1tkTghNOhsNTv89C6fbPnjliZ0TXgbFn5zE7JuRo4P1zmcVo5PqPH7gc4vDbmvwrAAD__yw3pRo">