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

git version: c7b537b

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)

[llvmbot]

@llvm/issue-subscribers-mlir-gpu

[joker-eph]

I suspect the input IR is invalid, you can't call a host function from the GPU code.

That said we should gracefully fail and not crash here.