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

    <tr>
        <th>Summary</th>
        <td>
            [MLIR] crashed on gpu::LaunchOp when using '-test-ir-visitors' pass
        </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: https://github.com/llvm/llvm-project/commit/9e91725fd4d4ee30e98ab2682f93b423590a4ade

Steps to reproduce:
```shell
mlir-opt test.mlir -test-ir-visitors
```
Test case:
```mlir
module {
  func.func @main(%arg0: tensor<4x4xf32>, %arg1: tensor<*xf32>, %arg2: tensor<*xf32>, %arg3: tensor<2x3xi32>, %arg4: tensor<2x3xi32>) attributes {pattern_driver_all_erased = true, pattern_driver_changed = false} {
    %c4 = arith.constant 4 : index
 %c0 = arith.constant 0 : index
    %0 = llvm.mlir.constant(0 : i32) : i32
 %1 = llvm.mlir.poison : vector<128xi32>
    %2 = llvm.mlir.constant(1 : index) : i64
    %3 = llvm.mlir.constant(dense<3.40282347E+38> : vector<128xf32>) : vector<128xf32>
    %4 = bufferization.to_memref %arg0 : tensor<4x4xf32> to memref<4x4xf32, strided<[?, ?], offset: ?>, #spirv.storage_class<StorageBuffer>>
    %alloc = memref.alloc() {alignment = 64 : i64} : memref<4xf32, #spirv.storage_class<StorageBuffer>>
    %5 = builtin.unrealized_conversion_cast %alloc : memref<4xf32, #spirv.storage_class<StorageBuffer>> to !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
    %6 = builtin.unrealized_conversion_cast %2 : i64 to index
    gpu.launch blocks(%arg5, %arg6, %arg7) in (%arg11 = %6, %arg12 = %6, %arg13 = %6) threads(%arg8, %arg9, %arg10) in (%arg14 = %6, %arg15 = %6, %arg16 = %6) {
      %7 = llvm.trunc %2 : i64 to i32
      %8 = llvm.insertelement %7, %1[%0 : i32] : vector<128xi32>
      %9 = llvm.shufflevector %8, %1 [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0] : vector<128xi32>
      %10 = llvm.icmp "sgt" %9, %1 : vector<128xi32>
      %11 = llvm.extractvalue %5[1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
      %12 = llvm.getelementptr %11[%2] : (!llvm.ptr, i64) -> !llvm.ptr, f32
      llvm.intr.masked.store %3, %12, %10 {alignment = 4 : i32} : vector<128xf32>, vector<128xi1> into !llvm.ptr
      gpu.terminator
    }
    return
  }
}

```
Crash trace:
```console
Block forward dominance post-order visits
Visiting block ^bb0mlir-opt: /home/workdir/llvm-project/llvm/include/llvm/ADT/ArrayRef.h:84: llvm::ArrayRef<mlir::BlockArgument>::ArrayRef(const T *, const T *) [T = mlir::BlockArgument]: Assertion `begin <= end' 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/workdir/llvm-project/build/bin/./mlir-opt test.mlir -test-ir-visitors
 #0 0x000055a5ee45b228 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x12d1228)
 #1 0x000055a5ee458d4e llvm::sys::RunSignalHandlers() (/home/workdir/llvm-project/build/bin/./mlir-opt+0x12ced4e)
 #2 0x000055a5ee45bc31 SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
 #3 0x00007fb094c7f520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
 #4 0x00007fb094cd39fc pthread_kill (/lib/x86_64-linux-gnu/libc.so.6+0x969fc)
 #5 0x00007fb094c7f476 gsignal (/lib/x86_64-linux-gnu/libc.so.6+0x42476)
 #6 0x00007fb094c657f3 abort (/lib/x86_64-linux-gnu/libc.so.6+0x287f3)
 #7 0x00007fb094c6571b (/lib/x86_64-linux-gnu/libc.so.6+0x2871b)
 #8 0x00007fb094c76e96 (/lib/x86_64-linux-gnu/libc.so.6+0x39e96)
 #9 0x000055a5eeb49ebf mlir::gpu::LaunchOp::getPrivateAttributions() (/home/workdir/llvm-project/build/bin/./mlir-opt+0x19bfebf)
#10 0x000055a5eeb49473 mlir::gpu::LaunchOp::verifyRegions() (/home/workdir/llvm-project/build/bin/./mlir-opt+0x19bf473)
#11 0x000055a5eec19055 mlir::Op<mlir::gpu::LaunchOp, mlir::OpTrait::OneRegion, mlir::OpTrait::VariadicResults, mlir::OpTrait::ZeroSuccessors, mlir::OpTrait::AtLeastNOperands<6u>::Impl, mlir::OpTrait::AttrSizedOperandSegments, mlir::OpTrait::OpInvariants, mlir::BytecodeOpInterface::Trait, mlir::OpTrait::AffineScope, mlir::OpTrait::AutomaticAllocationScope, mlir::InferIntRangeInterface::Trait, mlir::gpu::AsyncOpInterface::Trait, mlir::OpTrait::HasRecursiveMemoryEffects>::verifyRegionInvariants(mlir::Operation*) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x1a8f055)
#12 0x000055a5eec174f0 mlir::RegisteredOperationName::Model<mlir::gpu::LaunchOp>::verifyRegionInvariants(mlir::Operation*) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x1a8d4f0)
#13 0x000055a5f171439b (anonymous namespace)::OperationVerifier::verifyOpAndDominance(mlir::Operation&) Verifier.cpp:0:0
#14 0x000055a5f17137b6 mlir::verify(mlir::Operation*, bool) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x45897b6)
#15 0x000055a5f160a333 mlir::AsmState::AsmState(mlir::Operation*, mlir::OpPrintingFlags const&, llvm::DenseMap<mlir::Operation*, std::pair<unsigned int, unsigned int>, llvm::DenseMapInfo<mlir::Operation*, void>, llvm::detail::DenseMapPair<mlir::Operation*, std::pair<unsigned int, unsigned int>>>*, mlir::FallbackAsmResourceMap*) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x4480333)
#16 0x000055a5f1616a45 mlir::Block::printAsOperand(llvm::raw_ostream&, bool) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x448ca45)
#17 0x000055a5f1be2383 printBlock(mlir::Block*) TestVisitors.cpp:0:0
#18 0x000055a5f1be6f0c void llvm::function_ref<void (mlir::Block*)>::callback_fn<testNoSkipErasureCallbacks(mlir::Operation*)::$_1>(long, mlir::Block*) TestVisitors.cpp:0:0
#19 0x000055a5f1134f9e void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Block*)>, mlir::WalkOrder) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x3faaf9e)
#20 0x000055a5f1134f83 void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Block*)>, mlir::WalkOrder) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x3faaf83)
#21 0x000055a5f1be2086 (anonymous namespace)::TestIRVisitorsPass::runOnOperation() TestVisitors.cpp:0:0
#22 0x000055a5f1551cd3 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x43c7cd3)
#23 0x000055a5f1552572 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+0x43c8572)
#24 0x000055a5f1554d4e mlir::PassManager::run(mlir::Operation*) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x43cad4e)
#25 0x000055a5f154d2cb performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#26 0x000055a5f154cf23 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
#27 0x000055a5f15f8475 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+0x446e475)
#28 0x000055a5f1546b82 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+0x43bcb82)
#29 0x000055a5f1546e33 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x43bce33)
#30 0x000055a5f1547042 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (/home/workdir/llvm-project/build/bin/./mlir-opt+0x43bd042)
#31 0x000055a5ee43a9af main (/home/workdir/llvm-project/build/bin/./mlir-opt+0x12b09af)
#32 0x00007fb094c66d90 (/lib/x86_64-linux-gnu/libc.so.6+0x29d90)
#33 0x00007fb094c66e40 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x29e40)
#34 0x000055a5ee43a505 _start (/home/workdir/llvm-project/build/bin/./mlir-opt+0x12b0505)
```
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJzsWt1P5LqS_2vCi0UrsfP5wENoQBdp5gwCdFbal5ZjV9Je0k7Wdhi4f_3KTrrz0TQM53D23t17pZmQOOWqX5XL9RE31VpUEuDCiy696OqMdmbbqAv9E4Dys6LhrxePoA1izW4njEdytDWm1R7JPXzj4ZtKmG1XrFiz8_BNXT_v_5y3qvkvYMbDN8NUfJNBFiQ4KnnIQwDiQ5bSAscpLjNShJhEmU9DysHzc8_PHwy0GpkGKWhVwzsGVqife7Hf_9NbqGvPz3e1UOdNa5ABbVb2CZ3b23Ohzp-FFqZRejrP8_NeJaqXLO1ky7HhXQ3ISy49P0eo7CRb2QvyQn9HhfRw6uGIqsq3BjEgdaM8sg5fwpeSYI9ce3iNeopgRuHh_IgCf0hBZhT4hbyIBUV4kiJD1Bglis6Atgq11BhQcsOVeAa1oXW9AUU1cOSRK2RUB5brgoptqawGkpLWGrzk6mAdZDGw0L2kSpjtijVSGyoNsoM5EpLDiyW1dP5bdP6crmfZU1pfcmt6oPZwOtATbPXb3_YCgsWsthG6kY7oGZhx9glwurfPQRo-LS2YoNvLi8NxKjk9lYO0PrYmq9DHKSZhcu3hS5J65PoYUnlYslOvDiJ7axddWYISf6dGNHJlms0OdgrKwSd6I73hnHZP9ZSTYbxG2ijBgVsnjC49cuPci9x40ZW9a8pSgwsAdmzvfES3Qj2vtGkUrWDDaqq1R9YP_fOlA2iJp-BpXTfMKdCjWLkBt6cy61W0FpXcgTSOJg4PFrdOR_Ip9gH5H8MRDUYUtRFy1UkFtBZ_B75hjXwGpUUjN4xqM8X8Z6Vb23s4cM6ijeqYcXs-bY1y-67_Y5XFa0SVoq_WA9CLG-qN_tZoNlMs_oRieG9di2y6A6u2W9W0k2yLirphT_oQ9KIx7sTjbWJXT0h0IAv6rWjxTMIhfmuQTAYzZLYKKB_lpSNlNpnkHwkM3-IdvTUYzwSOkcyZLxn3s1Eu6i_N1EebPX060gupQRmooXdfHCWD0MBuqT6m9dEquvogJjnW2chab7uyrKGf4MTuWSMvuvTtw78v_778Iy6_5svBJKELtmuRh7GujIex8_TRm3-B1STLw4tRlJlnWnfgoroXXQZ7SH9tqO2xTCqHCvZ7vzWqB9rvezwCSveYpgAydO5KgvmrchJnhvBi1GpH9RNwl2ycwmRvOby_8Y-zaHgIPMkbizUWnHPDBxaUkJOcZZHtIdkMYUDthKR2zpB-kqvhToHplHQP_eBwnVfia0X1FtklXNbitoxqatsMXNr8g8pG_aSKI95YiZIBahttzhvFQSFX59sq_3d7I2TVJy3kRddF4e87hH4FbrbNDjx887NRT1yo44ZlaGOEZHXHYRzIrx7t1TrFPZSrrUfy1FXe7j3JPZLvX3pk7XoJN-jw56rq7HK4OmBGilNXMaJH5OHcLsLsMbPh_bEvl05wjK4siFzb3CNssRv7BVQ2NZK1nQeSezhBJRU18JXn53ffrvOHa6S7YicMoqjoKttlNcrY9PbpBk9o3YH28A2ikqPBbMhsATG3ugVlT26FV66to-wJ8W7X9uvtrzw_u1NNpegO0UEl_YsrZWscbv_anuxm5eGbX2wGbdHmI__F930_imgEEEYFxulkKfWr7m_ulJDGwX50borTkUjRn5tGGwV052FXYQhb-Gf9Rv8T-D186b8EmAfY1cUD5GABOeUhvAX5vpMPopK0_huVvAal99X116BiwEMYUeGlIRkJ0Ey-h1NnlzXSohKybDZm8PXnRvDBz_sZesVa6xq--98LIIOApCz8LGRJGWF_0KUWhYdvXtJ4E4fntZDdy3klu_4FW-lmFTvMIY6wPyIO5ww5yUqG2r7w3DyJuv4c9yzOSjZyj5ZwwyRGlXbqfRZ2mMQj43jOOI6SkiBa2H37KbY4TUoysk2O2AbFpxkGxcgwXRgghiz-HEOSQTZRPJt5WBFmUJSTcFi1XX_zzTUsP9phGMydEs_UQD58AxGN_NqtkBUlFGUP1O5Pfwk0TMjHQJ9BifL1Hqq_Al-YkBHfPICwIPOjaILPAlq_BxevZ9SPirqPgvZBQq_AOzS_UyUoF-wedFcb_Q7lf4JqHjrGQGsbsE8T5uYbUG1--9GCopLbtjvuDjn2dtfW70426sF2xsPsB6j67POOmu2tfLZqHJFdvhpgDQdLYUCVQ0VD8n72OyjKUkh4YE0L71F1ptlRI1he1w1zn3zemHIrS1C30txTWcHHOA4LnOtXyT6N_G9U3wPrlBbP8B12jXq9LktgNn1fH3v21HDplCMop86-3PkSx6dp6UfR6Ph44fhJWPoTtSxCbUANrmDh_EZ3gxm-NxzqDzbGP4HCPCz9UWEyUbgMkiAkmQvqVDbyddd0Gkm6A926iiZbIPvdaiFATXX60eaSX-2L71MauQ8q--mLVG5hhQtYJCniyTr0ot4x1xoVTVN_mdnCKM2SIh7NFs3wxT4lZBrAc717MNTA4uk9vNM3rpQUsrqpaaX7On8oG8ci7gqkhu90HomXTLXh_YuWWpJ1J90xDkdDmTV77lu7Ywm3smzeleKKs-VsDoaKes7prkfxdXj7f0sD3tC6ts1Ernf3oJtOMWeoL9xGYZj6hEwSZjz3hyCmYbRsxAbN7NrmekglH3QJX-vEYcpoOAl2yQx0AZikBDl8Pd6puw4jzoSPoM3vQ5P0xt5NF2zj0mfORybeUXaS2XXf9B_J3dtT8g5Bkw3ruimlR9a2XfuteXgS7bWiulOwHl6_G0b7QQ-Hm8C5Tlo3slpk6F9XNZupGpCwzKBXdWQ33Qg_af002wA3_deKW2NBNqqHdHp__HEDzjT8D1o__VDcNlxf5FukpLTM4OBb2D-yTEr-dS2TjqECB8td56fxRxnXuuHt_d4R79y5kYsYnfwhJ6ZIP_ZZjGfyoyhgnJxYlB_tY_OjteJyTlu3CoPUmUEdnuMsdiq_5ZLWr1ro71TSCtQY5xbx_cviHmEJ45MVIAsL4CjBn7bAnWihFnKZ1C3pQbH4y00ynWRF3UptlPsW1qgjrnOCXnL_YsB-RxVIY_P7vsj4whRJWBoleLR6uLB6yENYgN2r_6aX_SX1cEgYPXyksijnhV0UcswK1IIqG7XL2f7LwHs5-1DD6C1VwDetsZXMOOOhL0gqG9Jmtd2o6vdvt_frRhp4MUeL-r0W6kdrvlMh140sRTXhkaHJ2ze2frxQjpWYTCLnt6YSjNZ92306op6a4OH0oHsnxX93cKR73_0NJ91TW3EoaVebDYcabNl8as5Q8E3j_XIFTtULb5rwV9fyL9dnRHclaA3M9K2mej12jvccYCxv_Hl580-xMh-457wejco0TKZFtG5rYXLJ71TDQOs91n-I0_0f3BLvUkzfPRglZOWOnE4Mf1lDEkOYjA0JThfxKS5S_Lbn_3_bt1-VzQpWpJOcmy0NCrPPFHODDgUG29K-lMg_6xYfWeMLtYRJ8038hZaJH552m89o-b-lDvfDcdGWp3eEZrREO7r_DdOfP5cr_IyOhxGHc7n9uU7Ms08em-GMZ-M3xeU5XBxD6KPNxs7ZaEOV2Uy0-WUREE5EhEsbRX6EeuZfZ6XI38em_Y8NzvgF4RnJ6BlcBEkYpInvJ-nZ9iJJCx8IZTxNfJzwjKcFKfwwC9MkCkMenIkL7OPIJwEJcJhE6YqnLKOMx8B9HqUMvNCHHRX1yv2AolHVmTsvvwhIEBB8VtMCau1-fo2x80uMvejqTF04vYqu0l7o10IbPXIwwtTuJ9u2oPSiq_6gHThqJDr6No1-bkGiTgtZIQ8nx8fhOEEt1fqsU_XFHz_8H_R5vsD_EwAA__-Th7eL">