[Bug] ROCm Qwen3 TVM FFI encountered a Segfault

[Erxl]

🐛 Bug

To Reproduce

Steps to reproduce the behavior:

1. mlc_llm gen_config /mnt/mytmp/Qwen3-32B-abliterated/config.json --quantization q4f16_0 --model-type qwen3 --conv-template deepseek_r1_qwen --context-window-size 16384 -o qwen3-32b-abliterated-q4f16_0
2. mlc_llm convert_weight /mnt/mytmp/Qwen3-32B-abliterated/config.json --quantization q4f16_0 --model-type qwen3 -o qwen3-32b-abliterated-q4f16_0
3. mlc_llm serve qwen3-32b-abliterated-q4f16_0 --host 192.168.1.100 --port 8080 --mode server --device rocm --prefix-cache-mode radix

换成 q4f16_0 q4f16_1 q4f16_2 均有相同错误

Expected behavior

Environment

• Platform (e.g. WebGPU/Vulkan/IOS/Android/CUDA):ROCm 6.2.0

• Operating system (e.g. Ubuntu/Windows/MacOS/...):Ubuntu 22

• Device (e.g. iPhone 12 Pro, PC+RTX 3090, ...) 7900xtx

• How you installed MLC-LLM (conda, source):python -m pip install --pre -U -f https://mlc.ai/wheels mlc-llm-nightly-rocm62 mlc-ai-nightly-rocm62

• How you installed TVM-Unity (pip, source):

• Python version (e.g. 3.10): 3.11

• GPU driver version (if applicable):

• CUDA/cuDNN version (if applicable):

• TVM Unity Hash Tag (python -c "import tvm; print('\n'.join(f'{k}: {v}' for k, v in tvm.support.libinfo().items()))", applicable if you compile models):BUILD_STATIC_RUNTIME: OFF
  BUILD_DUMMY_LIBTVM: OFF
  COMPILER_RT_PATH: 3rdparty/compiler-rt
  CUDA_VERSION: NOT-FOUND
  DLPACK_PATH: 3rdparty/dlpack/include
  DMLC_PATH: 3rdparty/dmlc-core/include
  GIT_COMMIT_HASH: 2d2d2ea7763b3cf5ed42cda79315103cc82d2309
  GIT_COMMIT_TIME: 2025-07-09 10:06:40 -0400
  HIDE_PRIVATE_SYMBOLS: ON
  INDEX_DEFAULT_I64: ON
  INSTALL_DEV: OFF
  LLVM_VERSION: 18.0.0git
  MLIR_VERSION: NOT-FOUND
  PICOJSON_PATH: 3rdparty/picojson
  RANG_PATH: 3rdparty/rang/include
  ROCM_PATH: /opt/rocm
  SUMMARIZE: OFF
  TVM_CXX_COMPILER_PATH: /opt/rh/gcc-toolset-11/root/usr/bin/c++
  USE_ALTERNATIVE_LINKER: AUTO
  USE_ARM_COMPUTE_LIB_GRAPH_EXECUTOR: OFF
  USE_ARM_COMPUTE_LIB: OFF
  USE_BLAS: none
  USE_BNNS: OFF
  USE_BYODT_POSIT: OFF
  USE_COREML: OFF
  USE_CPP_RPC: OFF
  USE_CPP_RTVM:
  USE_CUBLAS: OFF
  USE_CUDA: OFF
  USE_NVTX: OFF
  USE_NCCL: OFF
  USE_MSCCL: OFF
  USE_CUDNN: OFF
  USE_CUSTOM_LOGGING: OFF
  USE_CUTLASS: OFF
  USE_AMX: OFF
  USE_DNNL: OFF
  USE_FALLBACK_STL_MAP: OFF
  USE_GTEST: AUTO
  USE_HEXAGON: OFF
  USE_HEXAGON_RPC: OFF
  USE_HEXAGON_SDK: /path/to/sdk
  USE_HEXAGON_GTEST: /path/to/hexagon/gtest
  USE_HEXAGON_EXTERNAL_LIBS: OFF
  USE_IOS_RPC: OFF
  USE_KHRONOS_SPIRV: OFF
  USE_LIBBACKTRACE: AUTO
  USE_LIBTORCH: OFF
  USE_LLVM: /opt/rocm/llvm/bin/llvm-config --ignore-libllvm --link-static
  USE_MLIR: OFF
  USE_METAL: OFF
  USE_MIOPEN: OFF
  USE_MKL: OFF
  USE_MRVL: OFF
  USE_MSVC_MT: OFF
  USE_NNPACK: OFF
  USE_OPENCL: OFF
  USE_OPENCL_ENABLE_HOST_PTR: OFF
  USE_OPENCL_EXTN_QCOM: NOT-FOUND
  USE_OPENCL_GTEST: /path/to/opencl/gtest
  USE_OPENMP: OFF
  USE_PAPI: OFF
  USE_RANDOM: ON
  TVM_DEBUG_WITH_ABI_CHANGE: OFF
  TVM_LOG_BEFORE_THROW: OFF
  USE_ROCBLAS: OFF
  USE_HIPBLAS: ON
  USE_ROCM: ON
  USE_RCCL: /opt/rocm/
  USE_RPC: ON
  USE_RTTI: ON
  USE_RUST_EXT: OFF
  USE_SORT: ON
  USE_SPIRV_KHR_INTEGER_DOT_PRODUCT: OFF
  USE_TENSORFLOW_PATH: none
  USE_TENSORRT_CODEGEN: OFF
  USE_TENSORRT_RUNTIME: OFF
  USE_TFLITE: OFF
  USE_THREADS: ON
  USE_THRUST: OFF
  USE_CURAND: OFF
  USE_VULKAN: ON
  USE_CLML: OFF
  TVM_CLML_VERSION:
  USE_CLML_GRAPH_EXECUTOR: OFF
  USE_UMA: OFF
  USE_MSC: OFF
  USE_CCACHE: AUTO
  USE_NVSHMEM: OFF
  USE_NNAPI_CODEGEN: OFF
  USE_NNAPI_RUNTIME: OFF
  BACKTRACE_ON_SEGFAULT: OFF

• Any other relevant information:

Additional context

[2025-07-18 12:06:11] INFO auto_device.py:79: Found device: rocm:0
[2025-07-18 12:06:11] INFO auto_device.py:79: Found device: rocm:1
[2025-07-18 12:06:11] INFO auto_device.py:79: Found device: rocm:2
[2025-07-18 12:06:11] INFO auto_device.py:79: Found device: rocm:3
[2025-07-18 12:06:11] INFO auto_device.py:79: Found device: rocm:4
[2025-07-18 12:06:11] INFO auto_device.py:79: Found device: rocm:5
[2025-07-18 12:06:11] INFO auto_device.py:79: Found device: rocm:6
[2025-07-18 12:06:11] INFO auto_device.py:79: Found device: rocm:7
[2025-07-18 12:06:11] INFO jit.py:43: MLC_JIT_POLICY = ON. Can be one of: ON, OFF, REDO, READONLY
[2025-07-18 12:06:11] INFO jit.py:118: Compiling using commands below:
[2025-07-18 12:06:11] INFO jit.py:119: /home/a/miniconda3/envs/mlcllm/bin/python -m mlc_llm compile qwen3-32b-abliterated-q4f16_0 --opt 'flashinfer=1;cublas_gemm=1;faster_transformer=0;cudagraph=1;cutlass=1;ipc_allreduce_strategy=NONE' --overrides '' --device rocm:0 --output /tmp/tmpby858nrj/lib.so
[2025-07-18 12:06:13] INFO auto_config.py:70: Found model configuration: qwen3-32b-abliterated-q4f16_0/mlc-chat-config.json
[2025-07-18 12:06:13] INFO auto_target.py:91: Detecting target device: rocm:0
[2025-07-18 12:06:14] INFO auto_target.py:93: Found target: {'kind': 'rocm', 'tag': '', 'keys': ['rocm', 'gpu'], 'mcpu': 'gfx1100', 'thread_warp_size': 32, 'max_threads_per_block': 1024, 'max_shared_memory_per_block': 65536, 'max_num_threads': 256, 'mtriple': 'amdgcn-amd-amdhsa-hcc'}
[2025-07-18 12:06:14] INFO auto_target.py:110: Found host LLVM triple: x86_64-unknown-linux-gnu
[2025-07-18 12:06:14] INFO auto_target.py:111: Found host LLVM CPU: znver3
[2025-07-18 12:06:14] INFO auto_config.py:154: Found model type: qwen3. Use --model-type to override.
Compiling with arguments:
--config Qwen3Config(hidden_act='silu', hidden_size=5120, intermediate_size=25600, attention_bias=False, num_attention_heads=64, num_hidden_layers=64, num_key_value_heads=8, rms_norm_eps=1e-06, rope_theta=1000000, vocab_size=151936, tie_word_embeddings=False, context_window_size=16384, prefill_chunk_size=2048, tensor_parallel_shards=1, head_dim=128, dtype='float32', max_batch_size=128, weight_block_size=None, kwargs={})
--quantization GroupQuantize(name='q4f16_0', kind='group-quant', group_size=32, quantize_dtype='int4', storage_dtype='uint32', model_dtype='float16', linear_weight_layout='KN', quantize_embedding=True, quantize_final_fc=True, num_elem_per_storage=8, num_storage_per_group=4, max_int_value=7, tensor_parallel_shards=0)
--model-type qwen3
--target {'kind': 'rocm', 'tag': '', 'keys': ['rocm', 'gpu'], 'host': {'kind': 'llvm', 'tag': '', 'keys': ['cpu'], 'mtriple': 'x86_64-unknown-linux-gnu', 'mcpu': 'znver3'}, 'libs': ['thrust', 'rocblas', 'miopen', 'hipblas'], 'mtriple': 'amdgcn-amd-amdhsa-hcc', 'max_num_threads': 256, 'max_shared_memory_per_block': 65536, 'max_threads_per_block': 1024, 'thread_warp_size': 32, 'mcpu': 'gfx1100'}
--opt flashinfer=0;cublas_gemm=0;faster_transformer=0;cudagraph=0;cutlass=0;ipc_allreduce_strategy=NONE
--system-lib-prefix ""
--output /tmp/tmpby858nrj/lib.so
--overrides context_window_size=None;sliding_window_size=None;prefill_chunk_size=None;attention_sink_size=None;max_batch_size=None;tensor_parallel_shards=None;pipeline_parallel_stages=None;disaggregation=None
[2025-07-18 12:06:14] INFO compile.py:140: Creating model from: Qwen3Config(hidden_act='silu', hidden_size=5120, intermediate_size=25600, attention_bias=False, num_attention_heads=64, num_hidden_layers=64, num_key_value_heads=8, rms_norm_eps=1e-06, rope_theta=1000000, vocab_size=151936, tie_word_embeddings=False, context_window_size=16384, prefill_chunk_size=2048, tensor_parallel_shards=1, head_dim=128, dtype='float32', max_batch_size=128, weight_block_size=None, kwargs={})
[2025-07-18 12:06:14] INFO compile.py:158: Exporting the model to TVM Unity compiler
[2025-07-18 12:06:17] INFO compile.py:164: Running optimizations using TVM Unity
[2025-07-18 12:06:17] INFO compile.py:186: Registering metadata: {'model_type': 'qwen3', 'quantization': 'q4f16_0', 'context_window_size': 16384, 'sliding_window_size': -1, 'attention_sink_size': -1, 'prefill_chunk_size': 2048, 'tensor_parallel_shards': 1, 'pipeline_parallel_stages': 1, 'disaggregation': False, 'kv_state_kind': 'kv_cache', 'max_batch_size': 128}
[2025-07-18 12:06:19] INFO pipeline.py:57: Running TVM Relax graph-level optimizations
[2025-07-18 12:06:23] INFO pipeline.py:57: Lowering to TVM TIR kernels
[12:06:25] /workspace/tvm/include/tvm/topi/transform.h:1219: Warning: Fast mode segfaults when there are out-of-bounds indices. Make sure input indices are in bound
[12:06:27] /workspace/tvm/include/tvm/topi/transform.h:1219: Warning: Fast mode segfaults when there are out-of-bounds indices. Make sure input indices are in bound
[2025-07-18 12:06:33] INFO pipeline.py:57: Running TVM TIR-level optimizations
[2025-07-18 12:07:14] INFO pipeline.py:57: Running TVM Dlight low-level optimizations
[2025-07-18 12:07:15] INFO pipeline.py:57: Lowering to VM bytecode
!!!!!!! TVM FFI encountered a Segfault !!!!!!!
File "", in pyx_pw_3tvm_3ffi_4core_8Function_1__call(object*, object* const*, long, object*)
File "", in pyx_f_3tvm_3ffi_4core_FuncCall(void*, object*, TVMFFIAny*, int*) [clone .constprop.0]
File "", in tvm::ffi::FunctionObj::SafeCall(void*, TVMFFIAny const*, int, TVMFFIAny*)
File "", in tvm::ffi::Function::FromTyped<tvm::transform::TVMFFIFuncReg8::{lambda(tvm::transform::Pass, tvm::ffi::RValueRef<tvm::IRModule, void>)#1}>(tvm::transform::TVMFFIFuncReg8::{lambda(tvm::transform::Pass, tvm::ffi::RValueRef<tvm::IRModule, void>)#1}, std::cxx11::basic_string<char, std::char_traits, std::allocator >)::{lambda(tvm::ffi::AnyView const*, int, tvm::ffi::Any*)#1}::operator()(tvm::ffi::AnyView const, int, tvm::ffi::Any)
File "", in tvm::transform::Pass::operator()(tvm::IRModule) const
File "", in tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
File "", in tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
File "", in tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
File "", in tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
File "", in std::Function_handler<tvm::tir::PrimFunc (tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext), tvm::tir::transform::ForceNarrowIndexToInt32()::{lambda(tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext)#1}>::M_invoke(std::Any_data const&, tvm::tir::PrimFunc&&, tvm::IRModule&&, tvm::transform::PassContext&&)
File "", in tvm::tir::ForceNarrowIndexToInt32(tvm::tir::PrimFunc)
File "", in tvm::tir::Int32DTypeNarrower::RewriteDataType(tvm::tir::PrimFunc)
File "", in tvm::tir::IndexDataTypeNormalizer::Rewrite(tvm::tir::PrimFunc)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#15}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::BlockRealizeNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#14}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::Int32DTypeNarrower::VisitStmt(tvm::tir::BlockNode const*)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::BlockNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#3}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::IfThenElseNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#3}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::IfThenElseNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#15}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::BlockRealizeNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#14}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::Int32DTypeNarrower::VisitStmt(tvm::tir::BlockNode const*)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::BlockNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::ForNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::ForNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::ForNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::ForNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::ForNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt(tvm::tir::ForNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#10}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::StmtMutator::VisitStmt(tvm::tir::SeqStmtNode const*)
File "", in tvm::ffi::Array<tvm::tir::Stmt, std::enable_if<storage_enabled_vtvm::tir::Stmt, void>::type> tvm::tir::StmtMutator::Internal::MutateArray<tvm::tir::Stmt, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, std::enable_if<storage_enabled_vtvm::tir::Stmt, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1})
File "", in tvm::ffi::ObjectPtrtvm::ffi::Object tvm::ffi::Array<tvm::tir::Stmt, void>::MapHelper<tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::ffi::ObjectPtrtvm::ffi::Object, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1})
File "", in tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#10}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::StmtMutator::VisitStmt_(tvm::tir::SeqStmtNode const*)
File "", in tvm::ffi::Array<tvm::tir::Stmt, std::enable_if<storage_enabled_vtvm::tir::Stmt, void>::type> tvm::tir::StmtMutator::Internal::MutateArray<tvm::tir::Stmt, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, std::enable_if<storage_enabled_vtvm::tir::Stmt, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1})
File "", in tvm::ffi::ObjectPtrtvm::ffi::Object tvm::ffi::Array<tvm::tir::Stmt, void>::MapHelper<tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::ffi::ObjectPtrtvm::ffi::Object, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1})
File "", in tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#15}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::BlockRealizeNode const*)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
File "", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)#14}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>)
File "", in tvm::tir::Int32DTypeNarrower::VisitStmt_(tvm::tir::BlockNode const*)
File "", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::BlockNode const*)
File "", in tvm::tir::IndexDataTypeRewriter::VisitBlockAnnotations(tvm::ffi::Map<tvm::ffi::String, tvm::ffi::Any, void> const&)
File "", in std::_Function_handler<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&), tvm::tir::IndexDataTypeRewriter::VisitBlockAnnotations(tvm::ffi::Map<tvm::ffi::String, tvm::ffi::Any, void> const&)::{lambda(tvm::ffi::ObjectRef const&)#1}>::_M_invoke(std::_Any_data const&, tvm::ffi::ObjectRef const&)
File "", in tvm::ffi::ObjectPtrtvm::ffi::Object tvm::ffi::Array<tvm::ffi::ObjectRef, void>::MapHelper<std::function<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&)>, tvm::ffi::ObjectRef>(tvm::ffi::ObjectPtrtvm::ffi::Object, std::function<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&)>)
File "", in std::_Function_handler<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&), tvm::tir::IndexDataTypeRewriter::VisitBlockAnnotations(tvm::ffi::Map<tvm::ffi::String, tvm::ffi::Any, void> const&)::{lambda(tvm::ffi::ObjectRef const&)#1}>::_M_invoke(std::_Any_data const&, tvm::ffi::ObjectRef const&)
File "", in tvm::ffi::ObjectPtrtvm::ffi::Object tvm::ffi::Array<tvm::ffi::ObjectRef, void>::MapHelper<std::function<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&)>, tvm::ffi::ObjectRef>(tvm::ffi::ObjectPtrtvm::ffi::Object, std::function<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&)>)
File "./signal/../sysdeps/unix/sysv/linux/x86_64/libc_sigaction.c", in 0x00007fc547d3851f
File "", in tvm::ffi::(anonymous namespace)::backtrace_handler(int)
File "", in tvm::ffi::(anonymous namespace)::Traceback()

Traceback (most recent call last):
File "/home/a/miniconda3/envs/mlcllm/bin/mlc_llm", line 8, in
sys.exit(main())
^^^^^^
File "/home/a/miniconda3/envs/mlcllm/lib/python3.11/site-packages/mlc_llm/main.py", line 50, in main
cli.main(sys.argv[2:])
File "/home/a/miniconda3/envs/mlcllm/lib/python3.11/site-packages/mlc_llm/cli/serve.py", line 208, in main
serve(
File "/home/a/miniconda3/envs/mlcllm/lib/python3.11/site-packages/mlc_llm/interface/serve.py", line 57, in serve
async_engine = engine.AsyncMLCEngine(
^^^^^^^^^^^^^^^^^^^^^^
File "/home/a/miniconda3/envs/mlcllm/lib/python3.11/site-packages/mlc_llm/serve/engine.py", line 895, in init
super().init(
File "/home/a/miniconda3/envs/mlcllm/lib/python3.11/site-packages/mlc_llm/serve/engine_base.py", line 590, in init
) = _process_model_args(models, device, engine_config)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/home/a/miniconda3/envs/mlcllm/lib/python3.11/site-packages/mlc_llm/serve/engine_base.py", line 171, in _process_model_args
model_args: List[Tuple[str, str]] = [_convert_model_info(model) for model in models]
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/home/a/miniconda3/envs/mlcllm/lib/python3.11/site-packages/mlc_llm/serve/engine_base.py", line 171, in
model_args: List[Tuple[str, str]] = [_convert_model_info(model) for model in models]
^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/home/a/miniconda3/envs/mlcllm/lib/python3.11/site-packages/mlc_llm/serve/engine_base.py", line 164, in _convert_model_info
model_lib = jit.jit(
^^^^^^^^
File "/home/a/miniconda3/envs/mlcllm/lib/python3.11/site-packages/mlc_llm/interface/jit.py", line 164, in jit
_run_jit(
File "/home/a/miniconda3/envs/mlcllm/lib/python3.11/site-packages/mlc_llm/interface/jit.py", line 124, in _run_jit
raise RuntimeError("Cannot find compilation output, compilation failed")
RuntimeError: Cannot find compilation output, compilation failed

[skyne98]

Same here! MI50 GPUs.

(mlcenv) root@fox-llm-workstation:/workspace# mlc_llm chat /models/MLC/Qwen3-32B-q4f16_1-MLC/
[2025-07-19 17:18:24] INFO auto_device.py:90: Not found device: cuda:0
[2025-07-19 17:18:29] INFO auto_device.py:79: Found device: rocm:0
[2025-07-19 17:18:29] INFO auto_device.py:79: Found device: rocm:1
[2025-07-19 17:18:29] INFO auto_device.py:79: Found device: rocm:2
[2025-07-19 17:18:29] INFO auto_device.py:79: Found device: rocm:3
[2025-07-19 17:18:31] INFO auto_device.py:90: Not found device: metal:0
[2025-07-19 17:18:34] INFO auto_device.py:79: Found device: vulkan:0
[2025-07-19 17:18:34] INFO auto_device.py:79: Found device: vulkan:1
[2025-07-19 17:18:34] INFO auto_device.py:79: Found device: vulkan:2
[2025-07-19 17:18:34] INFO auto_device.py:79: Found device: vulkan:3
[2025-07-19 17:18:34] INFO auto_device.py:79: Found device: vulkan:4
[2025-07-19 17:18:36] INFO auto_device.py:90: Not found device: opencl:0
[2025-07-19 17:18:39] INFO auto_device.py:79: Found device: cpu:0
[2025-07-19 17:18:39] INFO auto_device.py:35: Using device: rocm:0
[2025-07-19 17:18:39] INFO jit.py:43: MLC_JIT_POLICY = ON. Can be one of: ON, OFF, REDO, READONLY
[2025-07-19 17:18:39] INFO jit.py:118: Compiling using commands below:
[2025-07-19 17:18:39] INFO jit.py:119: /opt/mlcenv/bin/python3 -m mlc_llm compile /models/MLC/Qwen3-32B-q4f16_1-MLC --opt 'flashinfer=1;cublas_gemm=1;faster_transformer=0;cudagraph=1;cutlass=1;ipc_allreduce_strategy=NONE' --overrides '' --device rocm:0 --output /tmp/tmpktx6rjjg/lib.so
[2025-07-19 17:18:41] INFO auto_config.py:70: Found model configuration: /models/MLC/Qwen3-32B-q4f16_1-MLC/mlc-chat-config.json
[2025-07-19 17:18:41] INFO auto_target.py:91: Detecting target device: rocm:0
[2025-07-19 17:18:43] INFO auto_target.py:93: Found target: {'kind': 'rocm', 'tag': '', 'keys': ['rocm', 'gpu'], 'mcpu': 'gfx906', 'thread_warp_size': 64, 'max_threads_per_block': 1024, 'max_shared_memory_per_block': 65536, 'max_num_threads': 256, 'mtriple': 'amdgcn-amd-amdhsa-hcc'}
[2025-07-19 17:18:43] INFO auto_target.py:110: Found host LLVM triple: x86_64-unknown-linux-gnu
[2025-07-19 17:18:43] INFO auto_target.py:111: Found host LLVM CPU: znver2
[2025-07-19 17:18:43] INFO auto_config.py:154: Found model type: qwen3. Use `--model-type` to override.
Compiling with arguments:
  --config          Qwen3Config(hidden_act='silu', hidden_size=5120, intermediate_size=25600, attention_bias=False, num_attention_heads=64, num_hidden_layers=64, num_key_value_heads=8, rms_norm_eps=1e-06, rope_theta=1000000, vocab_size=151936, tie_word_embeddings=False, context_window_size=40960, prefill_chunk_size=2048, tensor_parallel_shards=1, head_dim=128, dtype='float32', max_batch_size=128, weight_block_size=None, kwargs={})
  --quantization    GroupQuantize(name='q4f16_1', kind='group-quant', group_size=32, quantize_dtype='int4', storage_dtype='uint32', model_dtype='float16', linear_weight_layout='NK', quantize_embedding=True, quantize_final_fc=True, num_elem_per_storage=8, num_storage_per_group=4, max_int_value=7, tensor_parallel_shards=0)
  --model-type      qwen3
  --target          {'kind': 'rocm', 'tag': '', 'keys': ['rocm', 'gpu'], 'host': {'kind': 'llvm', 'tag': '', 'keys': ['cpu'], 'mtriple': 'x86_64-unknown-linux-gnu', 'mcpu': 'znver2'}, 'libs': ['thrust', 'rocblas', 'miopen', 'hipblas'], 'mtriple': 'amdgcn-amd-amdhsa-hcc', 'max_num_threads': 256, 'max_shared_memory_per_block': 65536, 'max_threads_per_block': 1024, 'thread_warp_size': 64, 'mcpu': 'gfx906'}
  --opt             flashinfer=0;cublas_gemm=0;faster_transformer=0;cudagraph=0;cutlass=0;ipc_allreduce_strategy=NONE
  --system-lib-prefix ""
  --output          /tmp/tmpktx6rjjg/lib.so
  --overrides       context_window_size=None;sliding_window_size=None;prefill_chunk_size=None;attention_sink_size=None;max_batch_size=None;tensor_parallel_shards=None;pipeline_parallel_stages=None;disaggregation=None
[2025-07-19 17:18:43] INFO compile.py:140: Creating model from: Qwen3Config(hidden_act='silu', hidden_size=5120, intermediate_size=25600, attention_bias=False, num_attention_heads=64, num_hidden_layers=64, num_key_value_heads=8, rms_norm_eps=1e-06, rope_theta=1000000, vocab_size=151936, tie_word_embeddings=False, context_window_size=40960, prefill_chunk_size=2048, tensor_parallel_shards=1, head_dim=128, dtype='float32', max_batch_size=128, weight_block_size=None, kwargs={})
[2025-07-19 17:18:43] INFO compile.py:158: Exporting the model to TVM Unity compiler
[2025-07-19 17:18:46] INFO compile.py:164: Running optimizations using TVM Unity
[2025-07-19 17:18:46] INFO compile.py:186: Registering metadata: {'model_type': 'qwen3', 'quantization': 'q4f16_1', 'context_window_size': 40960, 'sliding_window_size': -1, 'attention_sink_size': -1, 'prefill_chunk_size': 2048, 'tensor_parallel_shards': 1, 'pipeline_parallel_stages': 1, 'disaggregation': False, 'kv_state_kind': 'kv_cache', 'max_batch_size': 128}
[2025-07-19 17:18:48] INFO pipeline.py:57: Running TVM Relax graph-level optimizations
[2025-07-19 17:18:54] INFO pipeline.py:57: Lowering to TVM TIR kernels
[17:18:56] /workspace/tvm/include/tvm/topi/transform.h:1219: Warning: Fast mode segfaults when there are out-of-bounds indices. Make sure input indices are in bound
[17:18:57] /workspace/tvm/include/tvm/topi/transform.h:1219: Warning: Fast mode segfaults when there are out-of-bounds indices. Make sure input indices are in bound
[2025-07-19 17:19:04] INFO pipeline.py:57: Running TVM TIR-level optimizations
[2025-07-19 17:19:52] INFO pipeline.py:57: Running TVM Dlight low-level optimizations
[2025-07-19 17:19:54] INFO pipeline.py:57: Lowering to VM bytecode
!!!!!!! TVM FFI encountered a Segfault !!!!!!!
  File "<unknown>", in __pyx_pw_3tvm_3ffi_4core_8Function_1__call__(_object*, _object* const*, long, _object*)
  File "<unknown>", in __pyx_f_3tvm_3ffi_4core_FuncCall(void*, _object*, TVMFFIAny*, int*) [clone .constprop.0]
  File "<unknown>", in tvm::ffi::FunctionObj::SafeCall(void*, TVMFFIAny const*, int, TVMFFIAny*)
  File "<unknown>", in tvm::ffi::Function::FromTyped<tvm::transform::__TVMFFIFuncReg8::{lambda(tvm::transform::Pass, tvm::ffi::RValueRef<tvm::IRModule, void>)#1}>(tvm::transform::__TVMFFIFuncReg8::{lambda(tvm::transform::Pass, tvm::ffi::RValueRef<tvm::IRModule, void>)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::ffi::AnyView const*, int, tvm::ffi::Any*)#1}::operator()(tvm::ffi::AnyView const, int, tvm::ffi::Any)
  File "<unknown>", in tvm::transform::Pass::operator()(tvm::IRModule) const
  File "<unknown>", in tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  File "<unknown>", in tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  File "<unknown>", in tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  File "<unknown>", in tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  File "<unknown>", in std::_Function_handler<tvm::tir::PrimFunc (tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext), tvm::tir::transform::ForceNarrowIndexToInt32()::{lambda(tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext)#1}>::_M_invoke(std::_Any_data const&, tvm::tir::PrimFunc&&, tvm::IRModule&&, tvm::transform::PassContext&&)
  File "<unknown>", in tvm::tir::ForceNarrowIndexToInt32(tvm::tir::PrimFunc)
  File "<unknown>", in tvm::tir::Int32DTypeNarrower::RewriteDataType(tvm::tir::PrimFunc)
  File "<unknown>", in tvm::tir::IndexDataTypeNormalizer::Rewrite(tvm::tir::PrimFunc)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#15}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::BlockRealizeNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#14}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::Int32DTypeNarrower::VisitStmt_(tvm::tir::BlockNode const*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::BlockNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#3}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::IfThenElseNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#3}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::IfThenElseNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#15}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::BlockRealizeNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#14}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::Int32DTypeNarrower::VisitStmt_(tvm::tir::BlockNode const*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::BlockNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#10}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::StmtMutator::VisitStmt_(tvm::tir::SeqStmtNode const*)
  File "<unknown>", in tvm::ffi::Array<tvm::tir::Stmt, std::enable_if<storage_enabled_v<tvm::tir::Stmt>, void>::type> tvm::tir::StmtMutator::Internal::MutateArray<tvm::tir::Stmt, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, std::enable_if<storage_enabled_v<tvm::tir::Stmt>, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1})
  File "<unknown>", in tvm::ffi::ObjectPtr<tvm::ffi::Object> tvm::ffi::Array<tvm::tir::Stmt, void>::MapHelper<tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::ffi::ObjectPtr<tvm::ffi::Object>, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1})
  File "<unknown>", in tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#10}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::StmtMutator::VisitStmt_(tvm::tir::SeqStmtNode const*)
  File "<unknown>", in tvm::ffi::Array<tvm::tir::Stmt, std::enable_if<storage_enabled_v<tvm::tir::Stmt>, void>::type> tvm::tir::StmtMutator::Internal::MutateArray<tvm::tir::Stmt, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, std::enable_if<storage_enabled_v<tvm::tir::Stmt>, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1})
  File "<unknown>", in tvm::ffi::ObjectPtr<tvm::ffi::Object> tvm::ffi::Array<tvm::tir::Stmt, void>::MapHelper<tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::ffi::ObjectPtr<tvm::ffi::Object>, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1})
  File "<unknown>", in tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::ForNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#15}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::BlockRealizeNode const*)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  File "<unknown>", in tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#14}::_FUN(tvm::ffi::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
  File "<unknown>", in tvm::tir::Int32DTypeNarrower::VisitStmt_(tvm::tir::BlockNode const*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitStmt_(tvm::tir::BlockNode const*)
  File "<unknown>", in tvm::tir::IndexDataTypeRewriter::VisitBlockAnnotations(tvm::ffi::Map<tvm::ffi::String, tvm::ffi::Any, void> const&)
  File "<unknown>", in std::_Function_handler<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&), tvm::tir::IndexDataTypeRewriter::VisitBlockAnnotations(tvm::ffi::Map<tvm::ffi::String, tvm::ffi::Any, void> const&)::{lambda(tvm::ffi::ObjectRef const&)#1}>::_M_invoke(std::_Any_data const&, tvm::ffi::ObjectRef const&)
  File "<unknown>", in tvm::ffi::ObjectPtr<tvm::ffi::Object> tvm::ffi::Array<tvm::ffi::ObjectRef, void>::MapHelper<std::function<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&)>, tvm::ffi::ObjectRef>(tvm::ffi::ObjectPtr<tvm::ffi::Object>, std::function<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&)>)
  File "<unknown>", in std::_Function_handler<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&), tvm::tir::IndexDataTypeRewriter::VisitBlockAnnotations(tvm::ffi::Map<tvm::ffi::String, tvm::ffi::Any, void> const&)::{lambda(tvm::ffi::ObjectRef const&)#1}>::_M_invoke(std::_Any_data const&, tvm::ffi::ObjectRef const&)
  File "<unknown>", in tvm::ffi::ObjectPtr<tvm::ffi::Object> tvm::ffi::Array<tvm::ffi::ObjectRef, void>::MapHelper<std::function<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&)>, tvm::ffi::ObjectRef>(tvm::ffi::ObjectPtr<tvm::ffi::Object>, std::function<tvm::ffi::ObjectRef (tvm::ffi::ObjectRef const&)>)
  File "./signal/../sysdeps/unix/sysv/linux/x86_64/libc_sigaction.c", in 0x00007f2cee1a432f
  File "<unknown>", in tvm::ffi::(anonymous namespace)::backtrace_handler(int)
  File "<unknown>", in tvm::ffi::(anonymous namespace)::Traceback()

Traceback (most recent call last):
  File "/opt/mlcenv/bin/mlc_llm", line 8, in <module>
    sys.exit(main())
             ^^^^^^
  File "/opt/mlcenv/lib/python3.12/site-packages/mlc_llm/__main__.py", line 46, in main
    cli.main(sys.argv[2:])
  File "/opt/mlcenv/lib/python3.12/site-packages/mlc_llm/cli/chat.py", line 36, in main
    chat(
  File "/opt/mlcenv/lib/python3.12/site-packages/mlc_llm/interface/chat.py", line 288, in chat
    JSONFFIEngine(
  File "/opt/mlcenv/lib/python3.12/site-packages/mlc_llm/json_ffi/engine.py", line 232, in __init__
    model_args = _process_model_args(models, device, engine_config)[0]
                 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/mlcenv/lib/python3.12/site-packages/mlc_llm/serve/engine_base.py", line 171, in _process_model_args
    model_args: List[Tuple[str, str]] = [_convert_model_info(model) for model in models]
                                         ^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/mlcenv/lib/python3.12/site-packages/mlc_llm/serve/engine_base.py", line 164, in _convert_model_info
    model_lib = jit.jit(
                ^^^^^^^^
  File "/opt/mlcenv/lib/python3.12/site-packages/mlc_llm/interface/jit.py", line 164, in jit
    _run_jit(
  File "/opt/mlcenv/lib/python3.12/site-packages/mlc_llm/interface/jit.py", line 124, in _run_jit
    raise RuntimeError("Cannot find compilation output, compilation failed")
RuntimeError: Cannot find compilation output, compilation failed

[delphiRo]

Are you crazty guys from mlc-ai? It seems you have broken the thing which 2 months ago just worked. I cann't now even start inference at all on all my servers which were equiped with gfx906 instinct. You also not to add the old wheels builds here: https://mlc.ai/wheels
So it is just impossible to revert back to worked one

[delphiRo]

I approved that install these WHLs: mlc_ai_nightly_rocm62-0.20.dev377-cp312-cp312-manylinux_2_28_x86_64.whl and mlc_llm_nightly_rocm62-0.20.dev51-cp312-cp312-manylinux_2_28_x86_64.whl under the rocm6.2 become previous worked well MLC to not working with many LLM models I have try to infere from HF.

[skyne98]

Are you crazty guys from mlc-ai?

Oh, come on, we don't need to be like that, it's an unpaid project work on by volunteers in the end! They are doing a great job otherwise!
But please, @junrushao, @MasterJH5574 maybe you guys have any insights?

[delphiRo]

Are you crazty guys from mlc-ai?

Oh, come on, we don't need to be like that, it's an unpaid project work on by volunteers in the end! They are doing a great job otherwise! But please, @junrushao, @MasterJH5574 maybe you guys have any insights?

Ok. Let's find the explanation about the fact why did you remove the previous versions of wheels? Where is the whl archive?

[delphiRo]

Yesterday I try to build the engine from sources and it seems that it keep the problem even on CPU inference too

[MasterJH5574]

Hi folks @Erxl @skyne98 @delphiRo, apologize for the convenience. We have fixed this issue and are building the wheels. Will report back when the wheels are ready so that you can upgrade. Thanks.

[MasterJH5574]

Hi folks, the wheel upgrade is done, and we have confirmed that with the latest wheel, the model compilation doesn't run into this error. So please upgrade to the latest package and try again. Sorry for the troubles and please let us know if anything goes wrong 🙏

[delphiRo]

Hi folks, the wheel upgrade is done, and we have confirmed that with the latest wheel, the model compilation doesn't run into this error. So please upgrade to the latest package and try again. Sorry for the troubles and please let us know if anything goes wrong 🙏

Thank you for quick attention. Can you please share the links to WHLs which were built with this error to be corrected.
It will also important to have a archive repo with previous builds. Do you have ones?
It will also be good to make a releases with WHLs attached for different ROCm versions combined with version of python.
So share the links and commands to update current installation
When I post this error, I tested on Python 3.12

[delphiRo]

It seems after bugfix it is working again

[delphiRo]

The bug still reproduce on nVidia using cuda-12.4 wheel. It seems that it wasn't build. So the attempt to install wheels using the command: python -m pip install --pre -U -f https://mlc.ai/wheels mlc-llm-nightly-cu124 mlc-ai-nightly-cu124
reproduce the same error: RuntimeError: Cannot find compilation output, compilation failed

[delphiRo]

@MasterJH5574

[delphiRo]

@MasterJH5574
The bug still reproduce on nVidia using cuda-12.4 wheel. It seems that it wasn't build. So the attempt to install wheels using the command: python -m pip install --pre -U -f https://mlc.ai/wheels mlc-llm-nightly-cu124 mlc-ai-nightly-cu124
reproduce the same error: RuntimeError: Cannot find compilation output, compilation failed

[delphiRo]

I checked and see that versions of wheels are similar by each one for ROCM (which renewed whl fixed this bug ) and CUDA.
mlc-ai-nightly-cu124 0.20.dev405
mlc-llm-nightly-cu124 0.20.dev56

mlc-ai-nightly-rocm62 0.20.dev405
mlc-llm-nightly-rocm62 0.20.dev56

But under the CUDA installation bug reproduced