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

    <tr>
        <th>Summary</th>
        <td>
            [MLIR] --gpu-kernel-outlining triggers 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>
          RiRi114
      </td>
    </tr>
</table>

<pre>
    Test on commit : 0f9720a6

steps to reproduce:
`
mlir-opt  temp.mlir --gpu-kernel-outlining 
`
test case:
```mlir
module {
  func.func @kernel_function() {
    // Placeholder for kernel functionality, operates on input data
    %data = memref.alloc() : memref<4x8xf32>
    %result = memref.alloc() : memref<4xf32>
    %c1 = arith.constant 1 : index

 gpu.launch blocks(%bx, %by, %bz) in (%grid_x = %c1, %grid_y = %c1, %grid_z = %c1)
    threads(%tx, %ty, %tz) in (%block_x = %c1, %block_y = %c1, %block_z = %c1) {
      %val = memref.load %data[%bx, %tx] : memref<4x8xf32>
 %sum_result = gpu.all_reduce add %val uniform {} : (f32) -> (f32)
 memref.store %sum_result, %result[%bx] : memref<4xf32>
 gpu.terminator
    }
    return
  }

  func.func @main() {
 %data_init = memref.alloc() : memref<4x8xf32>
    %result_init = memref.alloc() : memref<4xf32>
    // Kernel launch with specific grid and block sizes
    %c2 = arith.constant 2 : index
    gpu.launch blocks(%bx, %by, %bz) in (%grid_x = %c2, %grid_y = %c2, %grid_z = %c2)
    threads(%tx, %ty, %tz) in (%block_x = %c2, %block_y = %c2, %block_z = %c2) {
      // Kernel computation
      func.call @kernel_function() : () -> ()
 gpu.terminator
    }
    return
  }
}
```

crash trace:
```
mlir-opt: /home/fuzzing/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: /home/fuzzing/llvm-project/build/bin/mlir-opt temp.mlir --gpu-kernel-outlining
 #0 0x000000000113ce77 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x113ce77)
 #1 0x000000000113aa0e llvm::sys::RunSignalHandlers() (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x113aa0e)
 #2 0x000000000113d855 SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
 #3 0x0000718280502520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
 #4 0x00007182805569fc pthread_kill (/lib/x86_64-linux-gnu/libc.so.6+0x969fc)
 #5 0x0000718280502476 gsignal (/lib/x86_64-linux-gnu/libc.so.6+0x42476)
 #6 0x00007182804e87f3 abort (/lib/x86_64-linux-gnu/libc.so.6+0x287f3)
 #7 0x00007182804e871b (/lib/x86_64-linux-gnu/libc.so.6+0x2871b)
 #8 0x00007182804f9e96 (/lib/x86_64-linux-gnu/libc.so.6+0x39e96)
 #9 0x000000000193c5be 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 0x00000000012714f7 mlir::WalkResult mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<mlir::WalkResult (mlir::Operation*)>, mlir::WalkOrder) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x12714f7)
#11 0x000000000193ae42 (anonymous namespace)::GpuKernelOutliningPass::runOnOperation() KernelOutlining.cpp:0:0
#12 0x00000000042452f7 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x42452f7)
#13 0x0000000004245b61 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x4245b61)
#14 0x000000000424825b mlir::PassManager::run(mlir::Operation*) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x424825b)
#15 0x00000000042408af performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#16 0x0000000004240503 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&)::$_3>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer> >, llvm::raw_ostream&) MlirOptMain.cpp:0:0
#17 0x00000000042ec4e8 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/fuzzing/llvm-project/build/bin/mlir-opt+0x42ec4e8)
#18 0x000000000423a151 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer> >, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x423a151)
#19 0x000000000423a403 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x423a403)
#20 0x000000000423a612 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x423a612)
#21 0x000000000111a8d7 main (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x111a8d7)
#22 0x00007182804e9d90 (/lib/x86_64-linux-gnu/libc.so.6+0x29d90)
#23 0x00007182804e9e40 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x29e40)
#24 0x000000000111a435 _start (/home/fuzzing/llvm-project/build/bin/mlir-opt+0x111a435)
Aborted (core dumped)
```
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJzcWltT5LoR_jXiRcWULd8feJiFJdnKUlBAJY8u2ZY9CrLkSPIus78-JcnGF2a4nExOTs6ps2DL6u6vu79uNZ7BStGGE3IBoi8gujrDvd4JeXFP76nvh2eFqPYXj0RpKDgsRdtSDUGwhV6dJcjDMfC2wNsqTToFtYCSdFJUfUlAYNZB7AFv2zIqz0WnIdSk7TbmFp6fN11__kQkJ-xc9JpRTnkDX2S0MVliNVNk_zfCRqWoekYgSL4Abwth3fNyY35AEHpOaW5uNRUcoBSg7GUrhABdA3QN7xguyU6wikhYCwmdGBzFMKN6D9AlFB2RWBNlAkB512tYYY1fdEXmDoLgCraklaTeYMZEORoNtsMyCC7D5_S5DhAIvk7Ckqie6Q-Kr4VL3wpiSfVuUwquNOYa-laM8oo8u-TApus3DPe83MGCifJJWfVR8WzcMxf78eKXsUo5dBsaSav82dqw1oZddnl_ePnXfDkboOqdJLgarOrRqh6t6qVVC_GAWbf-2q5bXxqeZdtG6gdm8xgzgasxdYb2s1joZxBdvZU3gCLVt_kscSa8mLFcEkN8iKtqNNlzWgvZWjSJ0wpQanShDJ6D4Ot0a1QP6JQWkiztDOCGmxHxa6ATTANKE9lSjrWQI2eSq-FKEt1Lbm_c4usyajFd1c4QsZxy-lHKHmP8Z3SsNNji_Zsr1oHVP6neQdWRkta0hIaIEPPKkR0q-ouoWdGgQ0WDFkVjtp6oaNDhokGHiwadqmjQkaJBR4oGvS6aeZxL0Xa9xradjjssWUrM2Bs91xF-zvbBwU_T0_0cTwFH2FJitYNa4nJ9SswOHYfheidaAtB13f_6RXkD0DVjP9rzTop_klIPtwBdU16yviLTwkPfdUKaHZdYacqbzQ4E2yiOjd6KlEzvOwJQinstjJ9WLNiCYFtipQFKLcXgtRQtBCi2UYm-PArHfXOY2c3XDOuHfVsIdk_qrdbS5MkJLfYt90RXBsVWKSJN1CGIPaowCC4fhakYlP4dM2vRGI4hQMiCCi4f9-65eYhl07eEayhqSLnJNNa0YAQ6z3yAEEAJrDFlpNoAb3v3_ev24StUfWGGAQyLvjHnvpDaTAA7rTtlwFoCNVTv-mJTinYK6SryVKmeKEM2U7NDAqDeEejyW-DyyebY2H7QuHyCVd92LuPexpHxTopG4vbFGfXBtBc9ZZX5bXrd9cug8t6c4rph4EHv2Rv_8_2gJEkyY4DaK3dxJynXFvqjJStKp00S_8yF0pLg1qbpElKuXdLS3-wAQF-85wHQUHEABf4KLsYeOQT3vucPtOGY_RXzihGpxnI-ASJjc0KEVoiqNIrgwjZAqY3HJVS0obwWuQZoa-5_CFrZy2yQUJuyM7Tw7D9nIBgMJH6KUi_yUIS8wQ9GC4Cun9M4j8NzRnn_fN7w3j0oN0psYos5NCIT4nCpMIqzuoSda9b5EzXd8DPaMyM_aY_WcMMkho2y7n0WdpjEk-J4oTgkaVIHEBemZj-lFhnBSW3ySq1ffFqhX0wK06XCOiNZ_DmFgRGZFGYLhmVBGRVk1k__gdnTvRvlpkIYT7HcTSAHdwOUTuu39k8Ee-4ZPpreOhwBjJn2ldccBJdKV26ZcFwwklOj3PrmLzoCVbngJBdL203Xu4vvdii57UwNHIMwf3JPmkPLX8wEYFZHsD8w621pzg-M9AU0VbnCLTkSj6Xu2fpr7e9sNKfObEdFNKbMXf_E7MmFbKnhVlamUWT-Uvm1kD-xrL5pExkh3eSTYi74vhW9ghy3RHW2H2dO4i9d7wae27HR32E1NEXZ81s-i3L6IgVQwnBbVBigZIHtQMpW_r-T1GV4jjPuj-FY7Oabd2QP1ApKmTAnyXE-Z3AFf9XqzeG2PIxR4od1crjS32TXcf68k4LT9o9XPBhJfoJT2MXG9UgTOX_VIzEJ0akZ9YEELoaBEIURWiRwnqvb7lHcdsbitsKdTc5geBFXC2nd-I53yi3HbK-ousEcN8Q2jEIIZn733L4lq042nA3-TUkI1t4Xsf9p7-9oRxjlZMUus_XFqfjk4ZgLGVPfuNLSjuGm627f2uAsuwcD9jssCdffeC2g_dtpaAAniXgR-1PEw1XEUxQVK6Cj6wfZtW5Rp0BoMEwIoxVCL8U17IishWy3tsWod_6WeDm-1Q5LUuWdliC4nCQeRC9LctOY7jaGe0WQm-_f7i8F1-RZv0rmDaPyttM3mPJLwWvazHRkcPb0QLnHa-ciL5g10e-ioSVm7w1nxwTmo0vP6b968sr3G9IKuf_S17Uj98v-itS4ZzqvCCOaHJcJvsKhVR9PwbFh8GAMP5rM_75DE7wrihkptZkilZb71_R4iwLjLBDmwfKg_2Pk5h2GJkuGkjIk6cxz1TGqt7y6k6IkSo1g_ze8-38sizd3zJ89aEl5c0_qo8sn6b42wVP3TZfpD7Af-YeJ_6er21NE08Zrima2jmboBUejOYwV5Q67AWL7WT68F4kTeRh6wYuHyFt7GPvoJB7-Lq7EPppcWb0p9HFaJbDF48cN_9k7QKttsoVW75CyKvvkKzpkRCaFwVohCT2Y50YmVxpLnc88-bAJEs5MhOv4hEEEnfLTRCgMImdtWwipSQXtJwmS2FffpBqQjB93nFUXQZUFGT4jF34SelEUJn5ytrvwAlSFoR-VMQmxl5VVlaDCIxGKQxIUBTmjF8hDkYdQ7CUo8pMNybIijVOUxmkdBnUKQo-0mLKNwb8Rsjmz7-svfJSmQXLGcEGYst8dQMjyFCEQXZ3JC-tv0TcKhB6jSqtJg6aa2e8bmNESRFfHvgygJW0aItXv-wnHWS_ZxW__FGMIzI8L9O8AAAD__1PmBZs">