<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=https://github.com/llvm/llvm-project/issues/116344>116344</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 @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>
  }
}
```
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)
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzcWF1z46oZ_jX4hrFHAn1e-MIf6_bMbCaZ3dxrELyS6cqgAjon6a_vgBzbcuJssj1tp804tiTgeZ_n_QAEs1a2CmCJ0jVKtzM2uL02yydpmWqjWa3F8_IRrMNaYa4PB-kQXeG9c71FdIXIDpFdK91-qBdcHxDZdd3vLz_z3ui_AXeI7I5DyS5Lk6LOaJqIWJRlBDEvi5xmTQo5pwKKPKtB8AJQtEXRyjroLXYaG-iNFgMHbzQ0oSw6fsLtoZNmrnuHHVi38Hd4Pm_7YX7QYuhg7vS8loqZZ0S3jTYH5hDdIpK2_VBxfehlx5zUqjq2vWXCI2PO7G0OwRRG-Xq8x7gZFF_4L4ySyBmmbK8tVEo3uhOV0qri1lVS9YNDpEAkZaaNvIMdKKsNohvyRJ8aShD94nHroWnAyH8ErgvGOViLEd1iRIgBJhAhKN8iUuK5H3FCoU_kAuWFHcaIpNy6gMCMdPsF18o6phwWoLzQDUrXMSIbHKF0G4ZPyckAOgGMApzTli1OivGLNLI52czXlVfeMQ6V7sEwJWzFnDNVNUpK14iQhnUWvKww8nybblG-DWwQKd7wlu__muc7jjlrMOAGoz5K0JnhktCof_WeBR-hMX1OF9M84obZPXaG3U53RHaDNb7OZO0rkHNEdk9FVmXJvJNqeJq3akBkF8eI7BaL119S8W4Q4EsTkXX4jJ1173OLdYiuElp6KdVjjxHJrBOeDV1V98cuVc0sVPLQd4hufMmN7W0_jBebc109MtOC80G5AfNRBPolfALAXRUwCx9WlK4DT7rFH2VSbcf-_zojXxorvLIWjG_AiORuL61PtequkrYC1bIWxMgVkRw3THYgFmMwH75-WX3_gu1QH6TDDNdD6yc8bZyf-j4910prB7CI7DBTAh8Djd0e8JhYNeM_QnIdzX93jP_AYjj0p2yLFjj8PRjdGnbAzLTDAZTzPLBPn9N0i8hurw8-j_7Q5oeQPif_xCkYI0IjHD1FURSlgpUFoYTSusBBeoiIfbbjxYORygUxj6F0SHHuZNgflbbOADsgkvnoS-VC3pDiUg8i6-gpjnJC6xCqE4n4ikQEAt4i8W1Q32WrWPdXpkQHxr7k55t2sgZ8DV7YIddiBS_wBBGR4sh9fGwXvPeRi8L_CYgegXIgTZbETZqSCAcyxznjjdmikzVfWL3IAr2EpCSakEsmmCQpy4bj3u390lP9kF33OQNlVjZ8YiC9Jp3kGW5t0PlZ8kmeTbCzKbao84ZiVvsq-xQyKfKGTpDzacyihrPmRsBT4KyZDC4mg1lDSlb4wUxp9XzQg8WKHcD2IaPLMcn-0g93oaYe9TpU1AOzx_wzg7pX93618rV0TL5p57fzpbzkkaZlk4nmYioV4Jjsxuv7_lHf997mSrDeaXMyjUhxHhJIkZUvtvPDC2pXLSvFumcr7R1TrPVpvsG11p3_HVTYpIr3ijYhAJk4-9bXbHStCTj5tKYH2UMnFUy0jV1PVLM_XeTlIG_qN2WdCXOwNq9Qpx1Gy2PDkfsDM6Dcb6rROOzwAsJtPwInEz_GUz8ykmVwZf5F0JuZMPXHLbtNnGUwsUuu4lfHJcE9GL9CrLgHtD-Z5U9ru90zA6Lqnd-TnUd814PhcNcavys8uuYqmHdff_u20crBk3vl-LtOmvve3TGpNlo1sr3AKPFF6-ua8_rolT4mmvpiUfmqW8lZ9w3s0LmL5_6FIqyUBpqJmukARIqT_EHJvw_wSv4dHLR5Xoe3iom7BDRs6FwloAMHt8eM2zGywbeDUJ42bJx1nd98VI2abK4u_PTRcP7b9ZzZbSXrgLtv0ErrzPPr_HgvB8YuiCRVFICLTqv2P6vkncj8PEOTSYZmScIgvdBu-066lRIPRvtX0Re6_5W8-x-sind7XLZ9d0aq9hs0Nx_fmFRpSRmkk0k1vZp00qYhb6fz_1sx3lrvkqaZrnfZlYsywuhNFx2Xa75n48K8-mzofqbvJu-UMDrhnV_zTurbof0M718lmNRTxxZXO2WSihwfmFS3tstxKvIJQnm1i-ei_ORrDSlFGV1ikugaE5IIV5UfVlnHjKteKH7GCiRTK_G19jhJ8Yh_U32cnGt35V9VQPi-XBsIL-wgziZeDohmYklFSUs2g2Wc0zjLojiNZvtlnCVJSiJeZpzHlGdJFOdNXOY0KiJK6nQmlyQiSRzHaVQkhEYLoEmW1GVM4iyFjOQoieDAZLfwabLQpp2Fs4ZlHGc0SWYdq6Gzy_FcLKRMOBebmWU4oaiH1qIk6qR19ozgpOvC6bPfZ6F0--snBtgZ2bZg7MVJzI7JbjDwgROZ2WC65a-ftRxd8PuS_DMAAP__o-8Ebw">