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

    <tr>
        <th>Summary</th>
        <td>
            [MLIR] crashed in `-gpu-kernel-outlining` pass with error message: Assertion `isa<To>(Val) && "cast<Ty>() argument of incompatible type!"' failed
        </td>
    </tr>

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

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

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

<pre>
    test commit: [66da9f38f](https://github.com/llvm/llvm-project/commit/66da9f38f374e786b2f1c0ecdab0b651c94c4f27)

Step to reproduce:
```
mlir-opt test.mlir -gpu-kernel-outlining
```

test case:
```
#map = affine_map<(d0)[s0, s1] -> ((d0 - s0) ceildiv s1)>
module {
  func.func @main(%arg0: memref<89xi32, strided<[?], offset: ?>>, %arg1: memref<1xi32, strided<[?], offset: ?>>, %arg2: memref<59x3x9x39xf32, strided<[?, ?, ?, ?], offset: ?>>, %arg3: memref<1x1x9x1xf32, strided<[?, ?, ?, ?], offset: ?>>, %arg4: memref<i8, strided<[], offset: ?>>, %arg5: memref<i8, strided<[], offset: ?>>){
    %c3 = arith.constant 3 : index
 %c59 = arith.constant 59 : index
    %c39 = arith.constant 39 : index
 %c9 = arith.constant 9 : index
    %c1 = arith.constant 1 : index
    %c0 = arith.constant 0 : index
    gpu.launch blocks(%arg6, %arg7, %arg8) in (%arg12 = %c59, %arg13 = %c1, %arg14 = %c1) threads(%arg9, %arg10, %arg11) in (%arg15 = %c3, %arg16 = %c1, %arg17 = %c1) {
      scf.for %arg18 = %c0 to %c9 step %c1 {
        %c1_8 = arith.constant 1 : index
        %8 = affine.apply #map(%c39)[%c0, %c1]
          gpu.launch blocks(%arg19, %arg20, %arg21) in (%arg25 = %8, %arg26 = %c1_8, %arg27 = %c1_8) threads(%arg22, %arg23, %arg24) in (%arg28 = %c1_8, %arg29 = %c1_8, %arg30 = %c1_8) {
          gpu.terminator
        } {SCFToGPU_visited}
      }
 gpu.terminator
    } {SCFToGPU_visited}
    return
 }
}
```


Crash backtrace:
```
mlir-opt: /home/workdir/llvm-project/llvm/include/llvm/Support/Casting.h:566: decltype(auto) llvm::cast(const From &) [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: mlir-opt test.mlir -gpu-kernel-outlining
 #0 0x00005620458c1138 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/workdir/llvm-project/build/bin/mlir-opt+0x18bf138)
 #1 0x00005620458be7b5 llvm::sys::RunSignalHandlers() (/home/workdir/llvm-project/build/bin/mlir-opt+0x18bc7b5)
 #2 0x00005620458c2251 SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
 #3 0x00007fa8d0b38520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
 #4 0x00007fa8d0b8c9fc pthread_kill (/lib/x86_64-linux-gnu/libc.so.6+0x969fc)
 #5 0x00007fa8d0b38476 gsignal (/lib/x86_64-linux-gnu/libc.so.6+0x42476)
 #6 0x00007fa8d0b1e7f3 abort (/lib/x86_64-linux-gnu/libc.so.6+0x287f3)
 #7 0x00007fa8d0b1e71b (/lib/x86_64-linux-gnu/libc.so.6+0x2871b)
 #8 0x00007fa8d0b2fe96 (/lib/x86_64-linux-gnu/libc.so.6+0x39e96)
 #9 0x000056204621a54d mlir::WalkResult llvm::function_ref<mlir::WalkResult (mlir::Operation*)>::callback_fn<std::enable_if<!llvm::is_one_of<mlir::gpu::LaunchOp, mlir::Operation*, mlir::Region*, mlir::Block*>::value && std::is_same<mlir::WalkResult, mlir::WalkResult>::value, mlir::WalkResult>::type mlir::detail::walk<(mlir::WalkOrder)1, mlir::ForwardIterator, (anonymous namespace)::GpuKernelOutliningPass::runOnOperation()::'lambda'(mlir::gpu::LaunchOp), mlir::gpu::LaunchOp, mlir::WalkResult>(mlir::Operation*, (anonymous namespace)::GpuKernelOutliningPass::runOnOperation()::'lambda'(mlir::gpu::LaunchOp)&&)::'lambda'(mlir::Operation*)>(long, mlir::Operation*) KernelOutlining.cpp:0:0
#10 0x0000562045a01e57 mlir::WalkResult mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<mlir::WalkResult (mlir::Operation*)>, mlir::WalkOrder) (/home/workdir/llvm-project/build/bin/mlir-opt+0x19ffe57)
#11 0x0000562046219122 (anonymous namespace)::GpuKernelOutliningPass::runOnOperation() KernelOutlining.cpp:0:0
#12 0x000056204915e1a3 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (/home/workdir/llvm-project/build/bin/mlir-opt+0x515c1a3)
#13 0x000056204915ea42 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (/home/workdir/llvm-project/build/bin/mlir-opt+0x515ca42)
#14 0x000056204916121e mlir::PassManager::run(mlir::Operation*) (/home/workdir/llvm-project/build/bin/mlir-opt+0x515f21e)
#15 0x000056204915918b performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#16 0x0000562049158de4 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
#17 0x0000562049208ce5 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) (/home/workdir/llvm-project/build/bin/mlir-opt+0x5206ce5)
#18 0x0000562049151fd2 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (/home/workdir/llvm-project/build/bin/mlir-opt+0x514ffd2)
#19 0x0000562049152288 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (/home/workdir/llvm-project/build/bin/mlir-opt+0x5150288)
#20 0x00005620491524a2 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (/home/workdir/llvm-project/build/bin/mlir-opt+0x51504a2)
#21 0x000056204589fb07 main (/home/workdir/llvm-project/build/bin/mlir-opt+0x189db07)
#22 0x00007fa8d0b1fd90 (/lib/x86_64-linux-gnu/libc.so.6+0x29d90)
#23 0x00007fa8d0b1fe40 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x29e40)
#24 0x000056204589f5a5 _start (/home/workdir/llvm-project/build/bin/mlir-opt+0x189d5a5)
Aborted (core dumped)
```
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJzcWtty4yrWfhpyQ8WF0JGLXChOe_9df3cl1emZuXQhgWQmCDSAOsnbTyHZluRDDrM9e_ZMVXeswOJbRz6WcKi1olac34D4FsR3V7RzG21u7DPnlF0Vmr3eOG4dLHXTCAfCHIL4NkkYJVWYVSC-AzjbONdaEOYArwBe1cJtumJR6gbglZS_dh_XrdF_56UDeLUFw6s9UJhGPM2SAldBiXjJaIGKJA5KEpVRhVOACUA5QPmj4y10GhreGs26knu1KAcJ2v5DeSOFudatg97uhf8NXtdtd_3EjeLyWndOCiVUfbAMoHxwlNoToACHDW0hCO8grSqh-LqhLQiXAGcMeeviW4sAXkIbgPgOXoPwCwQ466fhNfRzBJZcSCZ-eRlMQPjFG6tZJzkE6S1AOYRVp8qF_wFBhBoqVI8RU1MjH_qGN4ZXIFxm5EWEuNfnjGCceVPiWxCu-owsoa4qy4d0hSuvKvzihwesYIYV_B4oPIOKyUv4Ql5C8lKdgfQLjz4-oCc8MDl4IS_BxbVEMy0iOwb_AEj8e0DIrhKgRyvDoeKMcJtFqZV1VDnoB3MoFOMvXtTLxeSUYD86kdyBnhQOyTHsScHToMEp0eCkKDolio5E67ZbSNqpcgMLqcsnu98LyRjsdHzM_BYTCu7FAtwrGuIzqf5wPxxMRqPpKIFuYzhlo84pAJo8B0da4z1QOJFLTipN50rH5ENoy2pRabOTzPaSyPPfkB3r2XAb_cnSbUbW2cdysl2QTdhtQdtWvsKB9QbXypAMPNfbsPWh9Gw3gXkra8EkhHgSQnwUQrwPYTYRm0RwPR1P5-PHqcN4Ij3JCY6OFGdndJDT4yE61D1PwxAOx00jFHXaTAOe3nnhx-Xqp_7t4S_rX8IKxxlIx3Bun09CvLvccNcZ1W_kfmz7c37gAZQvDbUbWNDyyRn61nHaExVebXTDAV49a_PEhDk-2rcHvlCl7BgfBx67ttXGSyypdULViw0I8zhJPC7jpXSvLQc4o53TPpD9sjAHYV5S6wDO-gKGK6MbCHDSxzq-_an7BHgbB-GVpO7xtSm0_MGr3DnjkzUsmsnNZeI7b0VuLTdOaAVBgoSlIFz-1D0rZ3-lstfoFScQYNwbFS5_vg7zfpKaumu4clBXUKhSNy11opAcDp4FwNdhCisqJGcLgPKHb1_yxy_QdkUjHKSw6Grf12jj_A7_dFMlrO24BXgFqWJwmwDoNhyW8xwv-j6Klk-QdU07ZBwthpp7MLo2tNk7Y_vD7OMdlacMBNELQgjFCUZRnJVBEGaTfNpXOzw8GKFcb8jPvvRwNgoZ-rzW1hlOmz7oSyiUG1KQfaAKi05I5j99B7XalzC-RS9BVlRBmA0dpTc3mJtb8LSIT5n7o1OPolZU_h9VTHJjt3m_hEVlWsSjRfgggBjHAZzpBjjr47GEVtRCVXrtAM7977-0YP0j2a6wi7L1SUb9_0FBuFWQVjRjqAizGKOtH1IUAK9esmSdRNdSqO7lulbdMFEurF4kvc0RjjEaLY7mgFlJqhK2AxGvn4SUn0MnCanKET0-NDdKE1jb3r3Pmh2lyQiczIEDnlYhpIXfgZ-CxVlahSNsegQbFJ8GDIoRMJsD4oqT5HOAIeFk4jiZVliCAxpHbMKOf6Py6Qe3nXSTjeBfSzw7rofG9qQ0wNk4ft9yQ_2KoR77FncgdCk9Ga0rBcKldWwY5ooWkq9F1b9WBaNiYdda8bWea63bbnj41ncb962v_nPKpzM_eH1q-NZ3K350Z-YvKju-o_y9lcKuLW34mQDMISfjM9B3pfyJMZFg3FEhh-dnKp-G9845wr1hnhZIMAdfafNMDfvqfDS0GVqXjCqtXhvdWahow23bsy8ZVvzWdv_f8_r9jtYfqN1SoOnUvZpENtuvAjiVtCkYBf5lPXszTWRu4juJnIfnfH39ORxLht7knbUndgbOpFb1WzVM4IH5B8Tuj7L50UtRwOP09L5-s7rO1887KbgsWxzVwa7IL3Dmkqri8e5eCYdBcMCIJMD40hX1gQTOjn4SxDyg4Zlc3bc_9X3rNeaMtn1ytopnce1NOiS78-yYKypfrbDfqaI17wmj0Fr6z071V4XsYq1YHMRlQMMxCeGh9zTCn_b-QbRcCsUPqsuL7p1KLh6O6SKv6quyzvQttGfd_C2BQfMwsbX9gRqu3FdVadi_92wJ4CIRpxEeIx7NI54EOOAHhu5cP1ldhxR1CQsrHPDRwvigJkiQFbDlptKmyXuKse-8OezPbruhhrN16wwIl-OKR92Zkn-vPbvtwn1QIN-_ff2x1MrxF3eUzO9SmPvWfadCLbWqRD3BIHAye2K7JwfOZYxHExL9pmtRUvleK3ZuAcDZ3vdOiX90_Mj377zR5vW2q6qhuPfyjFe0k27NuOSOn1-zv_w8n4Fznd_JEH40l_92f0br7gSVvHS-b7TOvB4Xx1sFsOsEojWaH_N_isy8U57prDwxykoeTxy3rRQuV-zB6JJbu7P1P1J0_4Vb4k2J6dyjM0LVP3h1dvgSxItRUvJ4JN7sgJuCiuHTVf-_tmcvcYpFVcUm5yw5CCbGWXY2mNuGotzQoXXIP1sO70XiMuc0wlm29xCjQw8jer5cPuPhH-JKRMdk4YMbQVIVKIUN3X1T8Pvu-ggr0PjegfHBXVHFyCev4jBhBI2A4SEgjxBcr_2atXXUuPXEkw-r4NFERXQYn5jGcAC_TIRiuuWhvNDGcQb7-3_D-wtrzraW7L6kuGI3ISMhoVf8JkhjnCZBmiVXm5uE4CoOeYR5SFCWpSSOcByzqMI8KEgaX4kbjHCM0iANoiDB8SIOWZoRyiIcpXFKGIgQb6iQC2__Qpv6qr9lvwkiEgbZlaQFl7b_0wmM-zrFGMR3V-am97foagsiJIV1dkRwwsn-zy18Uwniu-F6vn-PgCBBp2_WEwRbai18Fm4DuTHawIZbS2v-R39zcdUZefOvfzmxjdyvG_zPAAAA___yEz-q">