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

    <tr>
        <th>Summary</th>
        <td>
            [mlir] --gpu-kernel-outlining crashed with assertion failure "cast<Ty>() argument of incompatible type!"
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            new issue
      </td>
    </tr>

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

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

<pre>
    git version: c7b537bf0923df05254f9fa4722b298eb8f4790d

system: `Ubuntu 20.04.6 LTS (Focal Fossa)`

reproduced with: `mlir-opt --gpu-kernel-outlining a.mlir`

a.mlir:
```
func.func @func1() -> vector<7x23xi1> {
 %c6 = arith.constant 6 : index
  
  gpu.launch blocks(%arg0, %arg1, %arg2) in (%arg6 = %c6, %arg7 = %c6, %arg8 = %c6) threads(%arg3, %arg4, %arg5) in (%arg9 = %c6, %arg10 = %c6, %arg11 = %c6) {
 gpu.terminator
  } {SCFToGPU_visited}
  
  %cst_7 = arith.constant dense<false> : vector<7x23xi1>
  return %cst_7 : vector<7x23xi1>
}
func.func private @func2(%arg0: index, %arg1: index, %arg2: tensor<7x23xf32>) {
  %c6 = arith.constant 6 : index
  
 gpu.launch blocks(%arg3, %arg4, %arg5) in (%arg9 = %c6, %arg10 = %c6, %arg11 = %c6) threads(%arg6, %arg7, %arg8) in (%arg12 = %c6, %arg13 = %c6, %arg14 = %c6) {
    func.call @func1() : () -> vector<7x23xi1>
 gpu.terminator
  } {SCFToGPU_visited}
  
 return
}
```

trace:
```
mlir-opt: /data/bin/llvm-project/llvm/include/llvm/Support/Casting.h:566: decltype(auto) llvm::cast(const From&) [with To = mlir::FlatSymbolRefAttr; From = mlir::SymbolRefAttr]: Assertion `isa<To>(Val) && "cast<Ty>() argument of incompatible type!"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.      Program arguments: /data/bin/llvm-project/build/bin/mlir-opt --gpu-kernel-outlining reduced.mlir
 #0 0x0000558426f0fa2f llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/data/bin/llvm-project/build/bin/mlir-opt+0x14a8a2f)
 #1 0x0000558426f0ccb4 SignalHandler(int) Signals.cpp:0:0
 #2 0x00007faffa0b1420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420)
 #3 0x00007faff9b9c00b raise (/lib/x86_64-linux-gnu/libc.so.6+0x4300b)
 #4 0x00007faff9b7b859 abort (/lib/x86_64-linux-gnu/libc.so.6+0x22859)
 #5 0x00007faff9b7b729 (/lib/x86_64-linux-gnu/libc.so.6+0x22729)
 #6 0x00007faff9b8cfd6 (/lib/x86_64-linux-gnu/libc.so.6+0x33fd6)
 #7 0x0000558427554208 (anonymous namespace)::GpuKernelOutliningPass::createKernelModule(mlir::gpu::GPUFuncOp, mlir::SymbolTable const&) KernelOutlining.cpp:0:0
 #8 0x000055842755a805 _ZN4llvm12function_refIFN4mlir10WalkResultEPNS1_9OperationEEE11callback_fnIZNS1_6detail4walkILNS1_9WalkOrderE1ENS1_15ForwardIteratorEZN12_GLOBAL__N_122GpuKernelOutliningPass14runOnOperationEvEUlNS1_3gpu8LaunchOpEE_SF_S2_EENSt9enable_ifIXaantsrSt11disjunctionIJSt7is_sameIT2_S4_ESJ_ISK_PNS1_6RegionEESJ_ISK_PNS1_5BlockEEEE5valuesrSJ_IT3_S2_E5valueEST_E4typeES4_OT1_EUlS4_E_EES2_lS4_ KernelOutlining.cpp:0:0
 #9 0x0000558427558674 (anonymous namespace)::GpuKernelOutliningPass::runOnOperation() KernelOutlining.cpp:0:0
#10 0x00005584297d401e mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (/data/bin/llvm-project/build/bin/mlir-opt+0x3d6d01e)
#11 0x00005584297d44fa mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (/data/bin/llvm-project/build/bin/mlir-opt+0x3d6d4fa)
#12 0x00005584297d4d01 mlir::PassManager::run(mlir::Operation*) (/data/bin/llvm-project/build/bin/mlir-opt+0x3d6dd01)
#13 0x00005584297c51ab performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#14 0x00005584297c6664 processBuffer(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::MlirOptMainConfig const&, mlir::DialectRegistry&, llvm::ThreadPool*) MlirOptMain.cpp:0:0
#15 0x00005584297c6814 mlir::LogicalResult llvm::function_ref<mlir::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&)::'lambda'(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) MlirOptMain.cpp:0:0
#16 0x00005584298ba404 mlir::splitAndProcessBuffer(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>, llvm::raw_ostream&, bool, bool) (/data/bin/llvm-project/build/bin/mlir-opt+0x3e53404)
#17 0x00005584297bfc37 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (/data/bin/llvm-project/build/bin/mlir-opt+0x3d58c37)
#18 0x00005584297c6b3b mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (/data/bin/llvm-project/build/bin/mlir-opt+0x3d5fb3b)
#19 0x0000558426e49f2b main (/data/bin/llvm-project/build/bin/mlir-opt+0x13e2f2b)
#20 0x00007faff9b7d083 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x24083)
#21 0x0000558426ee141e _start (/data/bin/llvm-project/build/bin/mlir-opt+0x147a41e)
Aborted (core dumped)
```

</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzcWVtz27gV_jX0C8YaEARvD36QZWrrXV80kdJ28sIBSZDCGgJYAHTsf98BqAvJyIk3SbfbZmKRAg6-c8U5BxDRmjWC0isvvPbCmwvSma1UV0RI8bqTnb5frC8KWb1eNcyAZ6o0k8IL5qCMizCIixqmKKhqGKIQ12lNcIxQgdKEFkmN4xRWHrzx4Lz_1K_a0J1d7UXwY9EJ0wEEZxDPInC3WQMPJUtZEg6WUmviodSL4HC9oq2SVVfSCnxmZrsH2nGmLmVrwOVl03aXT1QJyi9lZzgTTDSAzCzFBGo_GOy_2tn-v_tad6Kc2Q_gYWifvocSD6Xg0gsy8ExLI5UXLOIXFLww34558XW_FHgoLCPgBTeAKGa2s1IKbYgwwA7OARMVfdmTgsOzabsZJ50ot6DgsnzSjl1IVAM9tAD9q396RVYWJsCRrGfoWJ-o4nODyXAwBWarKKlO_IITJT69hlN-6TloH54d9cccT5ayWhuqdkwQa8-DTeIbS7NeLDfyl9XH_JlpZmjlxTdTq1lIbfL4nLErKjT1gkVNuH1mzvZnHHfAUtR0Sgwhv0J-lOUUJ61iz8TQQ7yggQOPTh948ssxZMcMFfrEsQ6QZTmy2feE19vR9Z_39jS-huE5CMopRx-dBQ_OjuK3AgwA4DxUEs6nO9klj69v6p8Tpn1kTQJnkm_6T6NISd9KSYcs18u9rIghHloWTHhoyfnz7rJV8ndamv1XDy2ZKHlX0dPAumtbqSzFgmjDRDOzKTSMIotZ0ZKb15Z6KCGdkdYublkw94J5SbTxUOJiDSyV3HmoN3V4bRMx2EjngkNKDeZLTsz6dVdI_oHWc2OUF1y7hRO6MU14YyWZa02VYVLY5M408YLFRrqNkPydcMfWco-Ah5ATLFhsXvt5O0lU0-2oMEDWgIlS7lpiWMEp6LXzPYQ8FIOaME6rWW_c1V02X2dAd8WOGUBA0TVAUWssYCTYGtNqKy9aemjZMLPtillpjXCw7MQBTOuOag8tAREV2PsBmC0FpSJ6CwpSPjln79mvDSmfQNXt2qP34Qy4fyslG0V2R630O_xfdIxXx8lvlUdFXUXt6-GhhAUQwBcIIQzDBKOohjVB9SAg9KvuX1aKCePk37joRcmJSJHPudRGUdLHywIwYXr_Jd-lgoeu4YuPSUJQbbuDo7T-RNqyLDBYs0YQ_jciKk6Vh5I9835Yz8rWWhu6vyMQ2gPFNalrAgsfIwjyXFFtpKK5MnvZOSs8tHxJojzCl5yJ7uWyEV0_0fY5b6blDO4lxgiO5A2GbNIiLSEsgCJM03fglxY5csg4gLAYIeMxclwkYQpIId8l-QAZoSRMR8jhFDk-evLdmDEaY0ZjzKSsq-iPYQZBXUUjzHgYC3EYYgQTi3nsaIEgO6pbF61pH6m_tN1vbl88HrbFiuh9hJeKEkP76XtZddwG-SmDNW23x1h9XHaifGxtoE8z3IbYDOTy5z51TvidD8dkogtJYAjyTw_Ybhcf2XpmE2WuaH27fMCWqw__QfjTB6o7brLVw9rP08eWKmLpsizzfVsMbQLKa3H7yc5HFTWEcfyZ8KfbO7fCQjyqiqrMz-yAHy6l-kxUdWsslFTZpwcf5b_cPV7P7_L8IfcROm9DH6tOPIqTCM_ZR24hg6btkjvXlTy2WZavl_ka5Vn2sDYpFdZcOatv_0mIMFqtje9XTP--1_f217WJmc412dHbDcrXOM_Wv-a3699yp3H0gTZO3eFgeG17nyzLsvCZ8I5qZWc3gWPbD2XrTZ5hWyeyNc4fN36efeQWPM-yNcrt6_v8lk78lkQx_pEYHNtwX-i-KYlNjKM8nsYVhj4dRGfv-v79sd3Ix9YynVekdc3Qnvco4J1UaD6O8oFsk5m5IPxVM31PBGlsHl6AQkpun51wR8_qp5SFoIoq6NNjKrDK-1PlcU3-sPIr1lLOxHjX96RHnaKfbo3hIsvqVmijXAcg1ReoY4Kecz-xl31FFBXmVtTykIPmP8XguCYjg6OpwSvoT-Q8aH42tsaG-wkCVtAfCRiMBSxDnxSgpaqWajd3uUV_o4XRptq3QFuiaJW3trddnFasZadKet8oe6g4Jfyhu-7vbj8spDD0xXzhynvO1GNr7gkTCylq1oyKxmD2_HbHE_2iKMKgVbKkWl93de0aoXep1wn2r45-od493Un1eoAa0Fe0Jh03eUU5NfTtNUHmWvX36jyku2GE09LY5K6Net3Pn_hsXN-1cvtp_i5zhVNzJT4eMLyTDSsJ74vpgNGw7nrB4q0FHkr-PHO-7dTUEfXHuGPpH4k9sNNfMDzOu_2dW6Yn8VDMya6oiIfiv5JXUMKlaP5cw35FpG9vmGi0YZKCYDjcMLrlzMxFtZpknP-Gdv-Dm_SrFKdOoX_-aHWkYYAhHlXHeJwOi7oM4vMb7f8tTfxopxEmZRCPbJlMS0sRFG_act_ylVvSN3fzcSisjWKi-UDrb2v845rURVCMNBkdZiKK0xoVYEcO17XffY8TUFSjESsEJ3cMFUwCkOf2yJ9rQ5TJB3zfe-uAYRKMuIxviyj1sU9Bj_-jV1MxwYMzyLyQytAKuKtTRd0NH61OsoxveC-qq6BKg5Rc0Cs_Sv0kCBI_vNhehRUsY0ILSsIiqBKIYATDGJeI1BRhGl6wKwRRABMfwTSEfjgLcIxDGiZximNa0NLDkO4I4zOry0yq5sJdUV5FOArSC04KyrX7ARAhQT8DN-kh5IU3F-rK6V90jfYw5EwbfUIxzHD3y6ELyvDmrTtGd-25_9UOkOP9bk0Y7xT9_ovci07xq--_nnXq_zsAAP__LxjFLA">