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

    <tr>
        <th>Summary</th>
        <td>
            [MLIR]mlir::ArrayAttr::getValue() crashed with segmentation fault when using -gpu-module-to-binary
        </td>
    </tr>

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

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

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

<pre>
    Test commit: https://github.com/llvm/llvm-project/commit/6d2dfd37bd50b21ed90427052198bd1f06c761f8
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)
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzkW1lz46zS_jXkhooL0OqLXDiL3y9Vk5pUJnNuXVhCMt_IoANoxnl__SlAthYvk0wyybtUpWIbmubpp5umQTbVmpeCsQsQXYLo-ow2ZiXVxa1R2mjZVGdLmT9dPDJtYCbXa25AMIMrY2oNghkgc0DmJTerZjnJ5BqQeVV9376c10r-P8sMIPN2KJnHOcmLPEiWeYSWBLN8ikKSoIjgabrMcYHiLIlxkQJ0DdBMG1ZraCRUrFYybzJmJ3VdIEYQriuuzmVtoGHaTOwneF7Wzfla5k3Fzo08X3JB1RMEMWqHuf_GmUP1QF375z5CQII1rSEIriEtCi7YYk1rEFwBkuYIkCmILjUC5ApqDKJreA6CGwhI6rrhObR9U5gxXuX8u5UhUxDctFM51fjFunMEAZlBbV8uRzq9wZAao_iyMUxDkFyWdTPJpDCUC6YWXgQk17artRJCWDQim9h_EIRoTblwVkyHQoBEGfKAFTcrq1UbKgy0jTPIRc42Q2l8SBrvSQMS0aqS2cJrX7O1YsXENXU4aMVLsWbCOJk49Fri0NkSzNpRILjiAdkxYlWzTU1FvtArWrO-_mF7D4JbAod0Qi6M7BrxZjCVt9lrwW9hyEH9GV6Q57Hq5Ukvwia0rqunNqodoCjDPtScb220-bboulNiA6iijchWi36ELL4xJVhl104wG7bBZSWzbxpy4VdDRHaaF8O3U2hWitG8J7snMxCHVJV6K9gzuZVCe00j7x-i1wt2jjsoMz3g5eB3ejl8iZeDV3sZQsVMo8S2BSS7ThsAbWoZ-XmQHmyynxyKEAeBqhJtKWgJVyXeayG2BRDslNVGdT3B0Z5wT0u01xLvtSRH9aVHe6b7FhwwqrNqeiAXb0m5gkpmeTUpKmoWP6T6tiiVbOqF5n_6NAUIwYBcYUBIJz0erdh_84OjqVL0aZu5ZtBq2v4Lbvbyvwsin32XDa8MF5NGKEYr_ifLF5kU35nSXIpFRrXZ2rkNabs1j0LS6cMv04d-ps-nMucVu8lPGpGzAvadpY1qMuM30dZx7UvroS0rGG5ck1_9h1qngyXZW2RuIi40U-Y7rRq2C1wfAAREl2i7f_wmXHAELDwFLGiBBSC6xO8EDJAoOoVpu1pCEF2SdyYrPgUsaoFFILp0xL23K5NT6OIWXQyiy_Aj0KUfvAKnp9hJWnbSD1iBGJ1ClrbIpu-4BHfI8Clk0xYZRu-4Eu185BQqvK1RMP6odYhPJnuMtwDJRy1F7JO-rwJsAeD2_wnPJxu_jwZkPKKXkjXbmJ0WJ-w23Tgcj4lfsIe7CY7v4G7ipYIgulkuca989Z8DW3S6mwRYK5Y7Km0PGkNKjtW62FeaON4Vu1s3oVGpa7vSl1iWdHX3EaJ6qYltjKJZFzDJByx70ktIJTOsYvYsUhvlwXp6XAW-w5WOSt62gm0vHoZde8FFemmmkjTfYRhU0l5ZN9oK-UywpvaYY6o_d8oOhTAJjrOM37PA2QEKT9BMgjYK01eS3FoiFfOMbUvN3eL1Dd3o8aobHu0Gh7vu7e7dwbuwxxWDTCmp4JppTUsGuYZUw0JWlfyhuyu0Q4PRBN4rWSq6tuf3xnKkPR_zlVwzQOYbTkW5bgCZm5VibP_u0C7U3L5yAch8d-X3Ag1ratZNZfulpieuCbu7PwTRBiGEomiaxjnJCUa5I9Rfeugn7d_cKy7MF0Ozb4-KZgyQtBNS9MdCaqMYXQPiCjgujLsjeD7y7W0qXwIy_9LUtVS2-avgG9vAS0ErPeEiA8EsIfagjHFnBR5bERVLOILsdfwfFXnFFCDpd8lzQGa_E-fUHrN7MMkIJorT-BDZD40YoNXbO5dX49wCzOraQkMRCGZRBzAY85inERwT93udG-JgSFrYYkqKkEYZnuYhQXCxUMxli4UyPuXMG612c2zSeBGH5xUXzea8FI3vqP1F3ETLCQLkEm1wGBLU3XsBEkQ9AhKComWWRO7y3TtmZnPmzJj2Y8nMf2xubt3jrrB-ZcEbKSvdLnlA5lxkVZOzruH2AZD5pd_GZ7vbFuvEljSMHGm9O_144MkwyMLwhCFLVu7uw9_bilVrQ0xAMAuSzoRkYEKBl8RupCSlQoqntWw0FHTNdO3S0dQb4hPdF6a4q3RsuHZGl3Xj3_xx__XOCX6ubbY6IPG5KOwWz0X56dN_7h4VFbqihksH_FYYpgo37cHBj1SVzHyurbz2fLq8-KI105Lm4_ma08o3_3H_FZC5A1RItbZ8e1Me5aVL7YOlTcKOznRMJybRIfRmq_uPuvGqdaubMz3g83PNlCPFpdF_KI-BjSzSW1nTMY8Fi38Wljsqt9rvqW5zvWrEZ9Fj8qWJ_pX2pRYhDncP7HC_IEgwS6Npuuw5J2eG8mrr_0f5ubamzHJaG6l2Fg3CxNk6jpDjsTMTtHrSXN9RQUvmysSllO4-uBHuEW7e1hftcxkSLhL_VjqlUnkWfzWfDRj14N1LyxiyhCUdYXiPMBJBW1z0dvaiEZk1duGfh7jeFmRw40UyWlVLmn1bFAIEV38Dwm9sDShF-Wu1QLc7tA2z60dbCzx-mrdc0erGHn_0ZGWrgl6hYkkfVVJFEpLpy_h-m1j5NWvidGhNMAqhNGFtCHVeuvt0-3AlhWEb4xvYhmWNYTOnfhAy1v83rtcmWt8_8LgPkLj14TNZI1c9gl398ODEOrW3D18FN63s_my_lNcOVxE9MhyhJIlt8dAxGo4YxTgM_qJZ7A1TU4Smo-QUjXmIs_DFPNzzmlVcsNHmb0V35sVvTsw4hG6FNsqdq-2inZ0S8DP7jhb7PVVMmFtRyG0h8dKT38-4jx33ccd9POI-wGg6grwlYcc01frZRdYek29pTpqi3onCmpOMzUEB-Yk5Jwx5U6yxxRr3sKZDrEmSIQJrpmwp5BOi_skNhjZ5eyhfUcXyRW2PSlfdiC-yURm7KxUIbnql6dWRjD323V3F1efa3FEurqQoePkm5e1j7wR2Lmvb35uoZStIw1GSmI7ZSkMMayUzpvVlUxTuGPUsshrB_9uwPbLu2Fqqp62qnnzOCtpUZpGzihl2fIzdtG-ez2Bfri2JH1jJtVFPbX83z6O7E7h3KehVcfks8sMQW_LJjnyCxuTHQXrYzL-gEw6T-8ww76pJcqQkez-TjhP72gPEs8IiwmQUFngcFlHez7afZMkzWj0w3VTmeOF7bMBfhdufnn7-LeHfHabe05JTUf8-R7p-yJNByBcszRj7B4T2PyWrveKo2_dyMPZygPrHXF1X3MxEfj8qPz6CtL9hxJ2U6E5a_vXXa53uGY7d3r4aXnHD2bYeJ2Ofj07jSUgj_O-ocX532RAND2kkGjOdFdFRpttjdrai_kA9G4bPF6O4KB9YcaL5NE2_2_wIDS-vSTw2n2XHi-mXmP-xdsbxyM3DR2MkypMkhWvqfmDwUiBmH8L27W492xNj0F2vkbT_THaZYYzSAC4WFV9mC22oMosWzPMezGYTLSexeyRLQpQGu0eydq7p2NagYNDP0k7wuu9X-GkDxoKCdc-Cv7Byd5MEXe6AdrZMKgbzZl2zvAO5_SLIWX4R5NNgSs_YBU4QiYI0CtHZ6iJLSRIkKckDNEVhES-XcRYWCUM5yyKaBmf8giASYIJDlBBC0ARFURqjYhrEKYtQwkCI2JryauK-6yJVeca1bthFEoVBeFbRJau0-2EbId6tBETXZ-rCMbBsSg1CVHFtdKfAcFO5H8Pdfbp9ANH1s59yK6pXLIc_uFlBvU_TjxUTsNFclIe_e3LWqOrixb-rc-baMHUW_y8AAP__E9Wnpg">