hlo-translate failure: segv:mlir::stablehlo::foldConvert

[cxxhou]

hlo-translate fails to convert the attached StableHLO module to an HLO proto. run_hlo_module executes the same module successfully through its direct StableHLO frontend. Failure signature: segv:mlir::stablehlo::foldConvert.

The same module runs successfully on:

• run_hlo_module on Interpreter (StableHLO input)
• run_hlo_module on CPU (StableHLO input)
• run_hlo_module on GPU (StableHLO input)

Environment

• CPU: Intel(R) Core(TM) i9-14900HX
• GPU: NVIDIA GeForce RTX 4060 Laptop GPU
• CUDA Driver: 580.126.09
• XLA commit: 6f1f2d916dcef0428bfa7ea66765be539126c4f2
• StableHLO commit: e6f81ebd06b3509f2c7fa6175430aadbd4d724ca

IR

module @IrToHlo.32 attributes {mhlo.cross_program_prefetches = [], mhlo.input_output_alias = [], mhlo.is_dynamic = false, mhlo.use_auto_spmd_partitioning = false} {
  func.func @main(%arg0: tensor<1x2x1x1x1xcomplex<f32>>, %arg1: tensor<1x2x1x1x1xcomplex<f32>>, %arg2: tensor<1x2x1x1x1xcomplex<f32>>, %arg3: tensor<1x2x1x1x1xcomplex<f32>>, %arg4: tensor<1x2x1x1x1xcomplex<f32>>, %arg5: tensor<5x6xf32>, %arg6: tensor<2x3xi32>, %arg7: tensor<3xf32>) -> (tensor<1x2x2x5x1xcomplex<f32>>, tensor<2x5x1xi64>, tensor<1x2x5x1xi1>, tensor<5x6xf32>) {
    %cst = stablehlo.constant dense<(1.000000e+00,0.000000e+00)> : tensor<2x2x5x1xcomplex<f32>>
    %c = stablehlo.constant dense<0> : tensor<i64>
    %0 = stablehlo.concatenate %arg4, %arg3, %arg2, %arg1, %arg0, dim = 3 : (tensor<1x2x1x1x1xcomplex<f32>>, tensor<1x2x1x1x1xcomplex<f32>>, tensor<1x2x1x1x1xcomplex<f32>>, tensor<1x2x1x1x1xcomplex<f32>>, tensor<1x2x1x1x1xcomplex<f32>>) -> tensor<1x2x1x5x1xcomplex<f32>>
    %1 = stablehlo.reshape %0 : (tensor<1x2x1x5x1xcomplex<f32>>) -> tensor<1x2x5x1xcomplex<f32>>
    %2 = stablehlo.broadcast_in_dim %1, dims = [0, 1, 3, 4] : (tensor<1x2x5x1xcomplex<f32>>) -> tensor<1x2x2x5x1xcomplex<f32>>
    %3 = stablehlo.reshape %0 : (tensor<1x2x1x5x1xcomplex<f32>>) -> tensor<2x5x1xcomplex<f32>>
    %4 = stablehlo.broadcast_in_dim %3, dims = [1, 2, 3] : (tensor<2x5x1xcomplex<f32>>) -> tensor<2x2x5x1xcomplex<f32>>
    %5 = stablehlo.multiply %4, %cst : tensor<2x2x5x1xcomplex<f32>>
    %6 = stablehlo.reshape %5 : (tensor<2x2x5x1xcomplex<f32>>) -> tensor<1x2x2x5x1xcomplex<f32>>
    %7 = stablehlo.add %2, %6 : tensor<1x2x2x5x1xcomplex<f32>>
    %8 = stablehlo.real %1 : (tensor<1x2x5x1xcomplex<f32>>) -> tensor<1x2x5x1xf32>
    %9 = stablehlo.convert %8 : (tensor<1x2x5x1xf32>) -> tensor<1x2x5x1xi64>
    %10 = stablehlo.reduce(%9 init: %c) applies stablehlo.add across dimensions = [0] : (tensor<1x2x5x1xi64>, tensor<i64>) -> tensor<2x5x1xi64>
    %11 = stablehlo.minimum %9, %9 : tensor<1x2x5x1xi64>
    %12 = stablehlo.compare  GT, %11, %9 : (tensor<1x2x5x1xi64>, tensor<1x2x5x1xi64>) -> tensor<1x2x5x1xi1>
    %13 = stablehlo.convert %cst : (tensor<2x2x5x1xcomplex<f32>>) -> tensor<2x2x5x1xf32>
    %14 = stablehlo.reshape %13 : (tensor<2x2x5x1xf32>) -> tensor<20xf32>
    %15 = stablehlo.reshape %14 : (tensor<20xf32>) -> tensor<1x20xf32>
    %16 = stablehlo.broadcast_in_dim %15, dims = [0, 1] : (tensor<1x20xf32>) -> tensor<2x20xf32>
    %17 = stablehlo.reshape %16 : (tensor<2x20xf32>) -> tensor<40xf32>
    %18 = stablehlo.slice %17 [0:30] : (tensor<40xf32>) -> tensor<30xf32>
    %19 = stablehlo.reshape %18 : (tensor<30xf32>) -> tensor<5x6xf32>
    %20 = "stablehlo.scatter"(%19, %arg6, %arg7) <{indices_are_sorted = false, scatter_dimension_numbers = #stablehlo.scatter<inserted_window_dims = [0, 1], scatter_dims_to_operand_dims = [0, 1]>, unique_indices = false}> ({
    ^bb0(%arg8: tensor<f32>, %arg9: tensor<f32>):
      %21 = stablehlo.subtract %arg8, %arg9 : tensor<f32>
      stablehlo.return %21 : tensor<f32>
    }) : (tensor<5x6xf32>, tensor<2x3xi32>, tensor<3xf32>) -> tensor<5x6xf32>
    return %7, %10, %12, %20 : tensor<1x2x2x5x1xcomplex<f32>>, tensor<2x5x1xi64>, tensor<1x2x5x1xi1>, tensor<5x6xf32>
  }
}

Reproduction (hlo-translate — failed (rc=-11, SIGSEGV))

Command:

hlo-translate \
  --mlir-to-hlo \
  --emit-proto \
  43c16ccbd33bc947.mlir

Output:

PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace and instructions to reproduce the bug.
Stack dump:
0.	Program arguments: hlo-translate --mlir-to-hlo --emit-proto 43c16ccbd33bc947.mlir
 #0 0x00005afcc799e44b llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (hlo-translate+0xae5a44b)
 #1 0x00005afcc799bc03 llvm::sys::RunSignalHandlers() (hlo-translate+0xae57c03)
 #2 0x00005afcc799f18c SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
 #3 0x00007c7bfec45330 (/lib/x86_64-linux-gnu/libc.so.6+0x45330)
 #4 0x00007c7bfed89010 __memcpy_avx_unaligned_erms ./string/../sysdeps/x86_64/multiarch/memmove-vec-unaligned-erms.S:833:0
 #5 0x00005afcc769982e mlir::DenseElementsAttr::IntElementIterator::operator*() const (hlo-translate+0xab5582e)
 #6 0x00005afcc6e7db5a llvm::LogicalResult mlir::stablehlo::(anonymous namespace)::foldConvert<mlir::stablehlo::ConvertOp>(mlir::PatternRewriter&, mlir::stablehlo::ConvertOp, mlir::DenseTypedElementsAttr, mlir::RankedTensorType) StablehloAggressiveFolder.cpp:0:0
 #7 0x00005afcc6e7d474 mlir::stablehlo::(anonymous namespace)::FoldConvertOpPattern::matchAndRewrite(mlir::stablehlo::ConvertOp, mlir::PatternRewriter&) const StablehloAggressiveFolder.cpp:0:0
 #8 0x00005afcc7457a20 void llvm::function_ref<void ()>::callback_fn<mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<llvm::LogicalResult (mlir::Pattern const&)>)::$_0>(long) PatternApplicator.cpp:0:0
 #9 0x00005afcc7455211 mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<llvm::LogicalResult (mlir::Pattern const&)>) (hlo-translate+0xa911211)
#10 0x00005afcc742bd23 (anonymous namespace)::GreedyPatternRewriteDriver::processWorklist() GreedyPatternRewriteDriver.cpp:0:0
#11 0x00005afcc742b4d9 void llvm::function_ref<void ()>::callback_fn<(anonymous namespace)::RegionPatternRewriteDriver::simplify(bool*) &&::$_2>(long) GreedyPatternRewriteDriver.cpp:0:0
#12 0x00005afcc7429bf5 mlir::applyPatternsGreedily(mlir::Region&, mlir::FrozenRewritePatternSet const&, mlir::GreedyRewriteConfig, bool*) (hlo-translate+0xa8e5bf5)
#13 0x00005afcc6eb838d mlir::stablehlo::StablehloTargetIndependentOptimizationPass::runOnOperation() (hlo-translate+0xa37438d)
#14 0x00005afcc7609235 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (hlo-translate+0xaac5235)
#15 0x00005afcc76097d6 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (hlo-translate+0xaac57d6)
#16 0x00005afcc760aab8 mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool) (hlo-translate+0xaac6ab8)
#17 0x00005afcc760926b mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (hlo-translate+0xaac526b)
#18 0x00005afcc760bbd6 mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) (hlo-translate+0xaac7bd6)
#19 0x00005afcc760b9a5 mlir::PassManager::run(mlir::Operation*) (hlo-translate+0xaac79a5)
#20 0x00005afcbf3a2756 xla::(anonymous namespace)::ConvertStablehloToHloProtoInternal(mlir::ModuleOp, xla::HloProto*, bool, bool, bool) stablehlo.cc:0:0
#21 0x00005afcbf3a215c xla::(anonymous namespace)::ConvertStablehloToHloInternal(mlir::ModuleOp, bool, bool) stablehlo.cc:0:0
#22 0x00005afcbf3a20f2 xla::ConvertStablehloToHlo(mlir::ModuleOp) (hlo-translate+0x285e0f2)
#23 0x00005afcbf08bd9a MlirToHloTranslate(mlir::ModuleOp, llvm::raw_ostream&) hlo_translate.cc:0:0
#24 0x00005afcc6d7e900 std::_Function_handler<llvm::LogicalResult (std::shared_ptr<llvm::SourceMgr> const&, llvm::raw_ostream&, mlir::MLIRContext*), mlir::TranslateFromMLIRRegistration::TranslateFromMLIRRegistration(llvm::StringRef, llvm::StringRef, std::function<llvm::LogicalResult (mlir::Operation*, llvm::raw_ostream&)> const&, std::function<void (mlir::DialectRegistry&)> const&)::$_0>::_M_invoke(std::_Any_data const&, std::shared_ptr<llvm::SourceMgr> const&, llvm::raw_ostream&, mlir::MLIRContext*&&) Translation.cpp:0:0
#25 0x00005afcc6d7cbf6 llvm::LogicalResult llvm::function_ref<llvm::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>::callback_fn<mlir::mlirTranslateMain(int, char**, llvm::StringRef)::$_1>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) MlirTranslateMain.cpp:0:0
#26 0x00005afcc77239c8 llvm::LogicalResult llvm::function_ref<llvm::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>::callback_fn<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)::$_0>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef const&, llvm::raw_ostream&) ToolUtilities.cpp:0:0
#27 0x00005afcc772320b 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::MemoryBufferRef const&, llvm::raw_ostream&)>, llvm::raw_ostream&, llvm::StringRef, llvm::StringRef) (hlo-translate+0xabdf20b)
#28 0x00005afcc7723783 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) (hlo-translate+0xabdf783)
#29 0x00005afcc6d7a36c mlir::mlirTranslateMain(int, char**, llvm::StringRef) (hlo-translate+0xa23636c)
#30 0x00005afcbf08c0d5 main (hlo-translate+0x25480d5)
#31 0x00007c7bfec2a1ca __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#32 0x00007c7bfec2a28b call_init ./csu/../csu/libc-start.c:128:20
#33 0x00007c7bfec2a28b __libc_start_main ./csu/../csu/libc-start.c:347:5
#34 0x00005afcc7c5b99a _start (hlo-translate+0xb11799a)

Reference run (Interpreter (StableHLO input) — succeeds)

Command:

export XLA_FLAGS='--xla_backend_extra_options=xla_cpu_use_multi_output_fusion --xla_enable_enzyme_comms_opt=true --xla_gpu_async_dot=true --xla_gpu_enable_cudnn_layer_norm=true --xla_gpu_enable_dus_accumulator_zero_init_elimination=true --xla_gpu_enable_latency_hiding_scheduler=true --xla_gpu_enable_scatter_determinism_expander=true --xla_gpu_experimental_stream_annotation=true --xla_gpu_experimental_use_ragged_dot_fusion --xla_gpu_gemm_rewrite_size_threshold=0 --xla_gpu_multi_streamed_windowed_einsum=true --xla_recognize_reduction_optimization_level=1 --xla_reduce_window_rewrite_base_length=1'
run_hlo_module \
  --platform=Interpreter \
  --reference_platform= \
  --input_format=stablehlo \
  43c16ccbd33bc947.mlir

Output:

 ** Running 43c16ccbd33bc947.mlir**
Running HLO module with runner HloRunner...
... compiled and ran in 0.00456926s.
Skipping reference runner

Contact

• Email: ch395@njit.edu, zhihao.yao@njit.edu, benquike@gmail.com