hlo-translate failure: xla_builder.h:199 builder_ != nullptr

[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: xla_builder.h:199 builder_ != nullptr.

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 @DependentTupleElements_OneReadOnly attributes {mhlo.cross_program_prefetches = [], mhlo.input_output_alias = [], mhlo.is_dynamic = false, mhlo.use_auto_spmd_partitioning = false} {
  func.func private @DependentTupleElements_OneReadOnly.Body(%arg0: tensor<i32>, %arg1: tensor<8xf32>) -> (tensor<i32>, tensor<8xf32>) {
    %0 = stablehlo.convert %arg0 : (tensor<i32>) -> tensor<f32>
    %1 = stablehlo.broadcast_in_dim %0, dims = [] : (tensor<f32>) -> tensor<8xf32>
    %2 = stablehlo.add %arg1, %1 : tensor<8xf32>
    return %arg0, %2 : tensor<i32>, tensor<8xf32>
  }
  func.func private @DependentTupleElements_OneReadOnly.Condition(%arg0: tensor<i32>, %arg1: tensor<8xf32>) -> tensor<i1> {
    %c = stablehlo.constant dense<10> : tensor<i32>
    %0 = stablehlo.compare  LT, %arg0, %c : (tensor<i32>, tensor<i32>) -> tensor<i1>
    return %0 : tensor<i1>
  }
  func.func @main() -> (tensor<i32>, tensor<8xf32>) {
    %c = stablehlo.constant dense<0> : tensor<i32>
    %cst = stablehlo.constant dense<0.000000e+00> : tensor<8xf32>
    %0:2 = stablehlo.while(%iterArg = %c, %iterArg_0 = %cst) : tensor<i32>, tensor<8xf32>
    cond {
      %c_1 = stablehlo.constant dense<10> : tensor<i32>
      %1 = func.call @wrapped_27649(%iterArg) : (tensor<i32>) -> tensor<2x3xf32>
      %2 = stablehlo.compare  LT, %iterArg, %c_1 : (tensor<i32>, tensor<i32>) -> tensor<i1>
      %3 = stablehlo.convert %1 : (tensor<2x3xf32>) -> tensor<2x3xi1>
      %4 = stablehlo.reshape %3 : (tensor<2x3xi1>) -> tensor<6xi1>
      %5 = stablehlo.slice %4 [0:1] : (tensor<6xi1>) -> tensor<1xi1>
      %6 = stablehlo.reshape %5 : (tensor<1xi1>) -> tensor<i1>
      %7 = stablehlo.xor %2, %6 : tensor<i1>
      stablehlo.return %7 : tensor<i1>
    } do {
      %1 = stablehlo.convert %iterArg : (tensor<i32>) -> tensor<f32>
      %2 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor<f32>) -> tensor<8xf32>
      %3 = stablehlo.add %iterArg_0, %2 : tensor<8xf32>
      stablehlo.return %iterArg, %3 : tensor<i32>, tensor<8xf32>
    }
    return %0#0, %0#1 : tensor<i32>, tensor<8xf32>
  }
  func.func private @wrapped_27649(%arg0: tensor<i32>) -> tensor<2x3xf32> {
    %cst = stablehlo.constant dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf32>
    %cst_0 = stablehlo.constant dense<[[7.000000e+00, 8.000000e+00, 9.000000e+00], [1.000000e+01, 1.100000e+01, 1.200000e+01]]> : tensor<2x3xf32>
    %cst_1 = stablehlo.constant dense<[[1.300000e+01, 1.400000e+01, 1.500000e+01], [1.600000e+01, 1.700000e+01, 1.800000e+01]]> : tensor<2x3xf32>
    %0 = "stablehlo.case"(%arg0) ({
      stablehlo.return %cst : tensor<2x3xf32>
    }, {
      stablehlo.return %cst_0 : tensor<2x3xf32>
    }, {
      stablehlo.return %cst_1 : tensor<2x3xf32>
    }) : (tensor<i32>) -> tensor<2x3xf32>
    return %0 : tensor<2x3xf32>
  }
}

Reproduction (hlo-translate — failed (rc=-6, SIGABRT))

Command:

hlo-translate \
  --mlir-to-hlo \
  --emit-proto \
  7228dbd96aa861d3.mlir

Output:

WARNING: All log messages before absl::InitializeLog() is called are written to STDERR
F0000 00:00:1786849796.509514 3888157 xla_builder.h:199] Check failed: builder_ != nullptr
*** Check failure stack trace: ***
    @     0x580651b04814  absl::lts_20260107::log_internal::LogMessage::SendToLog()
    @     0x580651b04796  absl::lts_20260107::log_internal::LogMessage::Flush()
    @     0x580650287e4f  xla::While()
    @     0x5806492b7c9f  mlir::stablehlo::(anonymous namespace)::ExportXlaOp()
    @     0x58064928f497  mlir::(anonymous namespace)::ConvertToHloModule::Lower()
    @     0x5806492948c1  mlir::(anonymous namespace)::ConvertToHloModule::LowerBasicBlockAsFunction()
    @     0x580649291515  mlir::(anonymous namespace)::ConvertToHloModule::RunOnFunction()
    @     0x5806492893e6  mlir::ConvertMlirHloToHlo()
    @     0x580649282846  xla::(anonymous namespace)::ConvertStablehloToHloProtoInternal()
    @     0x58064928215c  xla::(anonymous namespace)::ConvertStablehloToHloInternal()
    @     0x5806492820f2  xla::ConvertStablehloToHlo()
    @     0x580648f6bd9a  MlirToHloTranslate()
    @     0x580650c5e900  std::_Function_handler<>::_M_invoke()
    @     0x580650c5cbf6  llvm::function_ref<>::callback_fn<>()
    @     0x5806516039c8  llvm::function_ref<>::callback_fn<>()
    @     0x58065160320b  mlir::splitAndProcessBuffer()
    @     0x580651603783  mlir::splitAndProcessBuffer()
    @     0x580650c5a36c  mlir::mlirTranslateMain()
    @     0x580648f6c0d5  main
    @     0x7c604d22a1ca  (unknown)
    @     0x7c604d22a28b  __libc_start_main
    @     0x580651b3b99a  _start
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 7228dbd96aa861d3.mlir
 #0 0x000058065187e44b llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (hlo-translate+0xae5a44b)
 #1 0x000058065187bc03 llvm::sys::RunSignalHandlers() (hlo-translate+0xae57c03)
 #2 0x000058065187f18c SignalHandler(int, siginfo_t*, void*) Signals.cpp:0:0
 #3 0x00007c604d245330 (/lib/x86_64-linux-gnu/libc.so.6+0x45330)
 #4 0x00007c604d29eb2c __pthread_kill_implementation ./nptl/pthread_kill.c:44:76
 #5 0x00007c604d29eb2c __pthread_kill_internal ./nptl/pthread_kill.c:78:10
 #6 0x00007c604d29eb2c pthread_kill ./nptl/pthread_kill.c:89:10
 #7 0x00007c604d24527e raise ./signal/../sysdeps/posix/raise.c:27:6
 #8 0x00007c604d2288ff abort ./stdlib/abort.c:81:7
 #9 0x0000580651b04586 (hlo-translate+0xb0e0586)
#10 0x0000580651b04874 absl::lts_20260107::log_internal::LogMessage::SendToLog() (hlo-translate+0xb0e0874)
#11 0x0000580651b04796 absl::lts_20260107::log_internal::LogMessage::Flush() (hlo-translate+0xb0e0796)
#12 0x0000580650287e4f xla::While(xla::XlaComputationId, xla::XlaComputationId, xla::XlaOp) (hlo-translate+0x9863e4f)
#13 0x00005806492b7c9f mlir::stablehlo::(anonymous namespace)::ExportXlaOp(mlir::stablehlo::WhileOp, (anonymous namespace)::OpLoweringContext) mlir_hlo_to_hlo.cc:0:0
#14 0x000058064928f497 mlir::(anonymous namespace)::ConvertToHloModule::Lower(mlir::Operation*, bool, llvm::ArrayRef<std::optional<xla::OpSharding>>, llvm::ArrayRef<mlir::Value>, xla::XlaBuilder*, llvm::DenseMap<mlir::Value, xla::XlaOp, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, xla::XlaOp>>*, xla::XlaOp*) mlir_hlo_to_hlo.cc:0:0
#15 0x00005806492948c1 mlir::(anonymous namespace)::ConvertToHloModule::LowerBasicBlockAsFunction(mlir::Block*, xla::XlaBuilder*, bool, bool, std::vector<bool, std::allocator<bool>> const&, llvm::ArrayRef<std::optional<xla::OpSharding>>, llvm::ArrayRef<std::optional<xla::FrontendAttributes>>, llvm::ArrayRef<std::optional<xla::OriginalValueProto>>, llvm::ArrayRef<std::optional<xla::OpSharding>>, xla::XlaComputationId&, llvm::ArrayRef<mlir::Value>, llvm::ArrayRef<mlir::Value>) mlir_hlo_to_hlo.cc:0:0
#16 0x0000580649291515 mlir::(anonymous namespace)::ConvertToHloModule::RunOnFunction(mlir::func::FuncOp) mlir_hlo_to_hlo.cc:0:0
#17 0x00005806492893e6 mlir::ConvertMlirHloToHlo(mlir::ModuleOp, xla::HloProto*, mlir::MlirToHloConversionOptions) (hlo-translate+0x28653e6)
#18 0x0000580649282846 xla::(anonymous namespace)::ConvertStablehloToHloProtoInternal(mlir::ModuleOp, xla::HloProto*, bool, bool, bool) stablehlo.cc:0:0
#19 0x000058064928215c xla::(anonymous namespace)::ConvertStablehloToHloInternal(mlir::ModuleOp, bool, bool) stablehlo.cc:0:0
#20 0x00005806492820f2 xla::ConvertStablehloToHlo(mlir::ModuleOp) (hlo-translate+0x285e0f2)
#21 0x0000580648f6bd9a MlirToHloTranslate(mlir::ModuleOp, llvm::raw_ostream&) hlo_translate.cc:0:0
#22 0x0000580650c5e900 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
#23 0x0000580650c5cbf6 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
#24 0x00005806516039c8 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
#25 0x000058065160320b 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)
#26 0x0000580651603783 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)
#27 0x0000580650c5a36c mlir::mlirTranslateMain(int, char**, llvm::StringRef) (hlo-translate+0xa23636c)
#28 0x0000580648f6c0d5 main (hlo-translate+0x25480d5)
#29 0x00007c604d22a1ca __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#30 0x00007c604d22a28b call_init ./csu/../csu/libc-start.c:128:20
#31 0x00007c604d22a28b __libc_start_main ./csu/../csu/libc-start.c:347:5
#32 0x0000580651b3b99a _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 \
  7228dbd96aa861d3.mlir

Output:

 ** Running 7228dbd96aa861d3.mlir**
Running HLO module with runner HloRunner...
... compiled and ran in 0.000649899s.
Skipping reference runner

Contact

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