[Bug] Error when trying to compile rwkv-raven-3b model.

🐛 Bug

Error when trying to compile rwkv-raven-3b model.

To Reproduce

Steps to reproduce the behavior:

python3 -m mlc_llm.build --hf-path=RWKV/rwkv-raven-3b --target iphone --quantization q8f16_1

Target configured: metal -keys=metal,gpu -libs=iphoneos -max_function_args=31 -max_num_threads=256 -max_shared_memory_per_block=32768 -max_threads_per_block=256 -thread_warp_size=1
Host CPU dection:
  Target triple: arm64-apple-darwin23.0.0
  Process triple: arm64-apple-darwin23.0.0
  Host CPU: apple-m1
Automatically using target for weight quantization: metal -keys=metal,gpu -max_function_args=31 -max_num_threads=256 -max_shared_memory_per_block=32768 -max_threads_per_block=1024 -thread_warp_size=32
Compilation error:
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:10:1: error: expected parameter declarator
;
^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:10:1: error: expected ')'
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:8:28: note: to match this '('
kernel void encode_kernel_1(  device half* max_abs_value [[ buffer(0) ]],
                           ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:8:61: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
kernel void encode_kernel_1(  device half* max_abs_value [[ buffer(0) ]],
                                                            ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:9:25: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device half* scale [[ buffer(1) ]],
                        ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:13:1: error: expected parameter declarator
;
^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:13:1: error: expected ')'
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:11:27: note: to match this '('
kernel void encode3_kernel(  device half* A [[ buffer(0) ]],
                          ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:11:48: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
kernel void encode3_kernel(  device half* A [[ buffer(0) ]],
                                               ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:12:33: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device half* max_abs_value [[ buffer(1) ]],
                                ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:17:1: error: expected parameter declarator
;
^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:17:1: error: expected ')'
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:14:28: note: to match this '('
kernel void encode_kernel_2(  device half* A [[ buffer(0) ]],
                           ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:14:49: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
kernel void encode_kernel_2(  device half* A [[ buffer(0) ]],
                                                ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:15:25: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device half* scale [[ buffer(1) ]],
                        ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:16:30: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device uint* w_gathered [[ buffer(2) ]],
                             ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:20:1: error: expected parameter declarator
;
^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:20:1: error: expected ')'
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:18:29: note: to match this '('
kernel void encode1_kernel_1(  device half* max_abs_value [[ buffer(0) ]],
                            ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:18:62: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
kernel void encode1_kernel_1(  device half* max_abs_value [[ buffer(0) ]],
                                                             ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:19:25: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device half* scale [[ buffer(1) ]],
                        ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:23:1: error: expected parameter declarator
;
^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:23:1: error: expected ')'
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:21:27: note: to match this '('
kernel void encode1_kernel(  device half* A [[ buffer(0) ]],
                          ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:21:48: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
kernel void encode1_kernel(  device half* A [[ buffer(0) ]],
                                               ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:22:33: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device half* max_abs_value [[ buffer(1) ]],
                                ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:26:1: error: expected parameter declarator
;
^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:26:1: error: expected ')'
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:24:26: note: to match this '('
kernel void encode_kernel(  device half* A [[ buffer(0) ]],
                         ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:24:47: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
kernel void encode_kernel(  device half* A [[ buffer(0) ]],
                                              ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:25:33: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device half* max_abs_value [[ buffer(1) ]],
                                ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:29:1: error: expected parameter declarator
;
^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:29:1: error: expected ')'
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:27:29: note: to match this '('
kernel void encode2_kernel_1(  device half* max_abs_value [[ buffer(0) ]],
                            ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:27:62: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
kernel void encode2_kernel_1(  device half* max_abs_value [[ buffer(0) ]],
                                                             ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:28:25: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device half* scale [[ buffer(1) ]],
                        ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:33:1: error: expected parameter declarator
;
^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:33:1: error: expected ')'
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:30:29: note: to match this '('
kernel void encode1_kernel_2(  device half* A [[ buffer(0) ]],
                            ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:30:50: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
kernel void encode1_kernel_2(  device half* A [[ buffer(0) ]],
                                                 ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:31:25: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device half* scale [[ buffer(1) ]],
                        ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:32:30: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device uint* w_gathered [[ buffer(2) ]],
                             ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:37:1: error: expected parameter declarator
;
^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:37:1: error: expected ')'
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:34:29: note: to match this '('
kernel void encode2_kernel_2(  device half* A [[ buffer(0) ]],
                            ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:34:50: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
kernel void encode2_kernel_2(  device half* A [[ buffer(0) ]],
                                                 ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:35:25: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device half* scale [[ buffer(1) ]],
                        ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:36:30: warning: 'buffer' attribute ignored on function declaration [-Wignored-attributes]
  device uint* w_gathered [[ buffer(2) ]],
                             ^
/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.metal:40:1: error: expected parameter declarator
;
^
fatal error: too many errors emitted, stopping now [-ferror-limit=]
21 warnings and 20 errors generated.
LLVM ERROR: Error opening '/var/folders/78/dqc2672d7nb8b5vp67c4n0n80000gn/T/tmpjmal_vsp/my_lib.air': No such file or directory!
Traceback (most recent call last):
  File "/opt/homebrew/anaconda3/envs/mlc-chat-venv/lib/python3.8/runpy.py", line 194, in _run_module_as_main
    return _run_code(code, main_globals, None,
  File "/opt/homebrew/anaconda3/envs/mlc-chat-venv/lib/python3.8/runpy.py", line 87, in _run_code
    exec(code, run_globals)
  File "/Users/scott/Developer/mlc-llm/mlc_llm/build.py", line 13, in <module>
    main()
  File "/Users/scott/Developer/mlc-llm/mlc_llm/build.py", line 10, in main
    core.build_model_from_args(parsed_args)
  File "/Users/scott/Developer/mlc-llm/mlc_llm/core.py", line 571, in build_model_from_args
    new_params = utils.convert_weights(param_manager, params, args)
  File "/Users/scott/Developer/mlc-llm/mlc_llm/utils.py", line 255, in convert_weights
    ex = relax.build(mod_transform, target=target)
  File "/opt/homebrew/anaconda3/envs/mlc-chat-venv/lib/python3.8/site-packages/tvm/relax/vm_build.py", line 338, in build
    return _vmlink(builder, target, tir_mod, ext_libs, params, system_lib=system_lib)
  File "/opt/homebrew/anaconda3/envs/mlc-chat-venv/lib/python3.8/site-packages/tvm/relax/vm_build.py", line 242, in _vmlink
    lib = tvm.build(
  File "/opt/homebrew/anaconda3/envs/mlc-chat-venv/lib/python3.8/site-packages/tvm/driver/build_module.py", line 281, in build
    rt_mod_host = _driver_ffi.tir_to_runtime(annotated_mods, target_host)
  File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__
  File "tvm/_ffi/_cython/./packed_func.pxi", line 262, in tvm._ffi._cy3.core.FuncCall
  File "tvm/_ffi/_cython/./packed_func.pxi", line 251, in tvm._ffi._cy3.core.FuncCall3
  File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL
tvm.error.InternalError: Traceback (most recent call last):
  File "/Users/catalyst/Workspace/mlc-ai-package-self-runner/_work/package/package/tvm/include/tvm/runtime/packed_func.h", line 838
InternalError: Check failed: type_code_ == kTVMStr (4 vs. 11) : expected str but got NULL

Expected behavior

Environment

Platform (e.g. WebGPU/Vulkan/IOS/Android/CUDA): iOS
Operating system (e.g. Ubuntu/Windows/MacOS/...): MacOS 14.0 Beta (23A5337a)
Device (e.g. iPhone 12 Pro, PC+RTX 3090, ...) M2
How you installed MLC-LLM (conda, source): source
How you installed TVM-Unity (pip, source): pip
Python version (e.g. 3.10): 3.8
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): `python -c "import tvm; print('\n'.join(f'{k}: {v}' for k, v in tvm.support.libinfo().items()))" ─

USE_GTEST: AUTO
SUMMARIZE: OFF
USE_IOS_RPC: OFF
USE_MSC: OFF
USE_ETHOSU: 
CUDA_VERSION: NOT-FOUND
USE_LIBBACKTRACE: AUTO
DLPACK_PATH: 3rdparty/dlpack/include
USE_TENSORRT_CODEGEN: OFF
USE_THRUST: OFF
USE_TARGET_ONNX: OFF
USE_AOT_EXECUTOR: ON
BUILD_DUMMY_LIBTVM: OFF
USE_CUDNN: OFF
USE_TENSORRT_RUNTIME: OFF
USE_ARM_COMPUTE_LIB_GRAPH_EXECUTOR: OFF
USE_CCACHE: AUTO
USE_ARM_COMPUTE_LIB: OFF
USE_CPP_RTVM: 
USE_OPENCL_GTEST: /path/to/opencl/gtest
USE_MKL: OFF
USE_PT_TVMDSOOP: OFF
MLIR_VERSION: NOT-FOUND
USE_CLML: OFF
USE_STACKVM_RUNTIME: OFF
USE_GRAPH_EXECUTOR_CUDA_GRAPH: OFF
ROCM_PATH: /opt/rocm
USE_DNNL: OFF
USE_VITIS_AI: OFF
USE_MLIR: OFF
USE_RCCL: OFF
USE_LLVM: llvm-config --link-static
USE_VERILATOR: OFF
USE_TF_TVMDSOOP: OFF
USE_THREADS: ON
USE_MSVC_MT: OFF
BACKTRACE_ON_SEGFAULT: OFF
USE_GRAPH_EXECUTOR: ON
USE_NCCL: OFF
USE_ROCBLAS: OFF
GIT_COMMIT_HASH: 6d1932ba16c44ce4cfb0e23dbadebe957916042a
USE_VULKAN: OFF
USE_RUST_EXT: OFF
USE_CUTLASS: OFF
USE_CPP_RPC: OFF
USE_HEXAGON: OFF
USE_CUSTOM_LOGGING: OFF
USE_UMA: OFF
USE_FALLBACK_STL_MAP: OFF
USE_SORT: ON
USE_RTTI: ON
GIT_COMMIT_TIME: 2023-09-06 14:13:53 -0400
USE_HEXAGON_SDK: /path/to/sdk
USE_BLAS: none
USE_ETHOSN: OFF
USE_LIBTORCH: OFF
USE_RANDOM: ON
USE_CUDA: OFF
USE_COREML: OFF
USE_AMX: OFF
BUILD_STATIC_RUNTIME: OFF
USE_CMSISNN: OFF
USE_KHRONOS_SPIRV: OFF
USE_CLML_GRAPH_EXECUTOR: OFF
USE_TFLITE: OFF
USE_HEXAGON_GTEST: /path/to/hexagon/gtest
PICOJSON_PATH: 3rdparty/picojson
USE_OPENCL_ENABLE_HOST_PTR: OFF
INSTALL_DEV: OFF
USE_PROFILER: ON
USE_NNPACK: OFF
LLVM_VERSION: 15.0.7
USE_OPENCL: OFF
COMPILER_RT_PATH: 3rdparty/compiler-rt
RANG_PATH: 3rdparty/rang/include
USE_SPIRV_KHR_INTEGER_DOT_PRODUCT: OFF
USE_OPENMP: OFF
USE_BNNS: OFF
USE_CUBLAS: OFF
USE_METAL: ON
USE_MICRO_STANDALONE_RUNTIME: OFF
USE_HEXAGON_EXTERNAL_LIBS: OFF
USE_ALTERNATIVE_LINKER: AUTO
USE_BYODT_POSIT: OFF
USE_HEXAGON_RPC: OFF
USE_MICRO: OFF
DMLC_PATH: 3rdparty/dmlc-core/include
INDEX_DEFAULT_I64: ON
USE_RELAY_DEBUG: OFF
USE_RPC: ON
USE_TENSORFLOW_PATH: none
TVM_CLML_VERSION: 
USE_MIOPEN: OFF
USE_ROCM: OFF
USE_PAPI: OFF
USE_CURAND: OFF
TVM_CXX_COMPILER_PATH: /Library/Developer/CommandLineTools/usr/bin/c++
HIDE_PRIVATE_SYMBOLS: ON

mlc-ai / mlc-llm

[Bug] Error when trying to compile rwkv-raven-3b model. #893

🐛 Bug

To Reproduce

Expected behavior

Environment