[MLIR]mlir::ArrayAttr::getValue() crashed with segmentation fault when using -gpu-module-to-binary #75434
Description
Test commit: 6d2dfd3
steps to reproduce:
mlir-opt test.mlir -gpu-module-to-binary
test case:
#map = affine_map<(d0)[s0, s1] -> ((d0 - s0) ceildiv s1)>
#map1 = affine_map<(d0)[s0, s1] -> (d0 * s0 + s1)>
module attributes {gpu.container_module} {
func.func @main() {
%c0 = arith.constant 0 : index
%c1 = arith.constant 1 : index
%alloc_0 = memref.alloc() {alignment = 64 : i64} : memref<i32>
%expand_shape = memref.expand_shape %alloc_0 [] : memref<i32> into memref<1xi32>
%alloc_1 = memref.alloc() {alignment = 64 : i64} : memref<1xi32>
%c1_2 = arith.constant 1 : index
%2 = affine.apply #map(%c1)[%c0, %c1]
gpu.launch_func @main_kernel::@main_kernel blocks in (%2, %c1_2, %c1_2) threads in (%c1_2, %c1_2, %c1_2) args(%c1 : index, %c0 : index, %expand_shape : memref<1xi32>, %alloc_1 : memref<1xi32>)
%alloc_3 = memref.alloc() {alignment = 64 : i64} : memref<1xi32>
%c1_4 = arith.constant 1 : index
%3 = affine.apply #map(%c1)[%c0, %c1]
return
}
gpu.module @main_kernel {
llvm.func @main_kernel(%arg0: i64, %arg1: i64, %arg2: !llvm.ptr, %arg3: !llvm.ptr, %arg4: i64, %arg5: i64, %arg6: i64, %arg7: !llvm.ptr, %arg8: !llvm.ptr, %arg9: i64, %arg10: i64, %arg11: i64) attributes {gpu.kernel, rocdl.flat_work_group_size = "1,1", rocdl.kernel, rocdl.reqd_work_group_size = array<i32: 1, 1, 1>} {
%0 = builtin.unrealized_conversion_cast %arg1 : i64 to index
%1 = builtin.unrealized_conversion_cast %arg0 : i64 to index
%2 = llvm.mlir.undef : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%3 = llvm.insertvalue %arg2, %2[0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%4 = llvm.insertvalue %arg3, %3[1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%5 = llvm.insertvalue %arg4, %4[2] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%6 = llvm.insertvalue %arg5, %5[3, 0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%7 = llvm.insertvalue %arg6, %6[4, 0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%8 = llvm.mlir.undef : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%9 = llvm.insertvalue %arg7, %8[0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%10 = llvm.insertvalue %arg8, %9[1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%11 = llvm.insertvalue %arg9, %10[2] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%12 = llvm.insertvalue %arg10, %11[3, 0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%13 = llvm.insertvalue %arg11, %12[4, 0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%14 = rocdl.workgroup.id.x : i32
%15 = llvm.sext %14 : i32 to i64
%16 = builtin.unrealized_conversion_cast %15 : i64 to index
llvm.br ^bb1
^bb1: // pred: ^bb0
%17 = affine.apply #map1(%16)[%1, %0]
%18 = builtin.unrealized_conversion_cast %17 : index to i64
%19 = llvm.extractvalue %7[1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%20 = llvm.getelementptr %19[%arg1] : (!llvm.ptr, i64) -> !llvm.ptr, i32
%21 = llvm.load %20 : !llvm.ptr -> i32
%22 = math.ctlz %21 : i32
%23 = llvm.extractvalue %13[1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
%24 = llvm.getelementptr %23[%18] : (!llvm.ptr, i64) -> !llvm.ptr, i32
llvm.store %22, %24 : i32, !llvm.ptr
llvm.return
}
}
}
The error message is as follows:
0. Program arguments: /home/xiangmu/three/llvm-project/build/bin/mlir-opt /home/xiangmu/three/llvm-project/matmul/tosa.mlir -gpu-module-to-binary
#0 0x000055986d2d210d llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) /home/xiangmu/three/llvm-project/llvm/lib/Support/Unix/Signals.inc:723:11
#1 0x000055986d2d25fb PrintStackTraceSignalHandler(void*) /home/xiangmu/three/llvm-project/llvm/lib/Support/Unix/Signals.inc:798:1
#2 0x000055986d2d0686 llvm::sys::RunSignalHandlers() /home/xiangmu/three/llvm-project/llvm/lib/Support/Signals.cpp:105:5
#3 0x000055986d2d2d85 SignalHandler(int) /home/xiangmu/three/llvm-project/llvm/lib/Support/Unix/Signals.inc:413:1
#4 0x00007f4a5c19d420 __restore_rt (/usr/lib/x86_64-linux-gnu/libpthread.so.0+0x14420)
#5 0x000055987205bc75 mlir::ArrayAttr::getValue() const /home/xiangmu/three/llvm-project/build/tools/mlir/include/mlir/IR/BuiltinAttributes.cpp.inc:103:10
#6 0x000055986d43c445 mlir::ArrayAttr::begin() const /home/xiangmu/three/llvm-project/build/tools/mlir/include/mlir/IR/BuiltinAttributes.h.inc:62:37
#7 0x000055986df1b222 (anonymous namespace)::moduleSerializer(mlir::gpu::GPUModuleOp, mlir::gpu::OffloadingLLVMTranslationAttrInterface, mlir::gpu::TargetOptions const&) /home/xiangmu/three/llvm-project/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp:105:24
#8 0x000055986df1b125 mlir::gpu::transformGpuModulesToBinaries(mlir::Operation*, mlir::gpu::OffloadingLLVMTranslationAttrInterface, mlir::gpu::TargetOptions const&) /home/xiangmu/three/llvm-project/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp:139:20
#9 0x000055986df1bfe6 (anonymous namespace)::GpuModuleToBinaryPass::runOnOperation() /home/xiangmu/three/llvm-project/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp:89:14
#10 0x0000559871e8598b mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7::operator()() const /home/xiangmu/three/llvm-project/mlir/lib/Pass/Pass.cpp:0:17
#11 0x0000559871e85925 void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7>(long) /home/xiangmu/three/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:5
#12 0x000055986d2f7429 llvm::function_ref<void ()>::operator()() const /home/xiangmu/three/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:5
#13 0x0000559871e887e5 void mlir::MLIRContext::executeAction<mlir::PassExecutionAction, mlir::Pass&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pass&) /home/xiangmu/three/llvm-project/mlir/include/mlir/IR/MLIRContext.h:276:3
#14 0x0000559871e81143 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /home/xiangmu/three/llvm-project/mlir/lib/Pass/Pass.cpp:509:17
#15 0x0000559871e816c4 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /home/xiangmu/three/llvm-project/mlir/lib/Pass/Pass.cpp:569:16
#16 0x0000559871e83109 mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) /home/xiangmu/three/llvm-project/mlir/lib/Pass/Pass.cpp:880:10
#17 0x0000559871e83032 mlir::PassManager::run(mlir::Operation*) /home/xiangmu/three/llvm-project/mlir/lib/Pass/Pass.cpp:860:60
#18 0x0000559871e77c02 performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) /home/xiangmu/three/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:384:17
#19 0x0000559871e77841 processBuffer(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::MlirOptMainConfig const&, mlir::DialectRegistry&, llvm::ThreadPool*) /home/xiangmu/three/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:441:12
#20 0x0000559871e77638 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&)::$_2::operator()(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) const /home/xiangmu/three/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:512:12
#21 0x0000559871e775d2 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&)::$_2>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) /home/xiangmu/three/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#22 0x0000559871fe8cee llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>::operator()(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) const /home/xiangmu/three/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12
#23 0x0000559871fe830d 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) /home/xiangmu/three/llvm-project/mlir/lib/Support/ToolUtilities.cpp:28:12
#24 0x0000559871e74a51 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) /home/xiangmu/three/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:515:10
#25 0x0000559871e74cf5 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) /home/xiangmu/three/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:550:14
#26 0x0000559871e74ec8 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) /home/xiangmu/three/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:566:10
#27 0x000055986d25d778 main /home/xiangmu/three/llvm-project/mlir/tools/mlir-opt/mlir-opt.cpp:284:33
#28 0x00007f4a5bc11083 __libc_start_main (/usr/lib/x86_64-linux-gnu/libc.so.6+0x24083)
#29 0x000055986d25d3fe _start (/home/xiangmu/three/llvm-project/build/bin/mlir-opt+0x23ee3fe)
Segmentation fault (core dumped)