Qwen2-72B-Instruct MultiGPU 8xP100

[alphaarea]

🐛 Bug

The just released Qwen2 has the same architecture as the previous Qwen1.5, so theoretically it should be able to run directly. In fact, the model was quantized and compiled without errors.

Starting a chat with mlc_llm chat, the model loads at [2/885] and says Aborted (core dumped). Other than that, there are no other error messages

To Reproduce

MODEL_PATH='/home/alphaarea/models/Qwen2-72B-Instruct'
MLC_QUANT='q4f16_1'
MLC_DEV='cuda'
MODEL_ARCH='qwen2'
MODEL_TEMP='chatml'
MODEL_NAME=${MODEL_PATH##*/}
MODEL_OUTPUT=$MODEL_PATH'-'$MLC_QUANT
MODEL_LIB=$MODEL_NAME'-'$MLC_QUANT'-'$MLC_DEV'.so'

mlc_llm convert_weight --quantization $MLC_QUANT --model-type $MODEL_ARCH --device $MLC_DEV --output $MODEL_OUTPUT $MODEL_PATH
mlc_llm gen_config --quantization $MLC_QUANT --model-type $MODEL_ARCH --conv-template $MODEL_TEMP --context-window-size 16384 --tensor-parallel-shards 8 --max-batch-size 1 --output $MODEL_OUTPUT $MODEL_PATH
mlc_llm compile --device $MLC_DEV --output $MODEL_OUTPUT/$MODEL_LIB $MODEL_OUTPUT/mlc-chat-config.json

mlc_llm chat --model-lib $MODEL_OUTPUT/$MODEL_LIB $MODEL_OUTPUT

Expected behavior

Starting a conversation as normal as in Qwen1.5

Environment

• Platform (e.g. WebGPU/Vulkan/IOS/Android/CUDA): CUDA
• Operating system (e.g. Ubuntu/Windows/MacOS/...): Ubuntu 22.04
• Device (e.g. iPhone 12 Pro, PC+RTX 3090, ...): Tesla P100 16G x8
• How you installed MLC-LLM (conda, source): python -m pip install --pre -U -f https://mlc.ai/wheels mlc-llm-nightly-cu121 mlc-ai-nightly-cu121 (date: 06/07/2024)
• How you installed TVM-Unity (pip, source): same as above
• Python version (e.g. 3.10): 3.11
• GPU driver version (if applicable): NVIDIA Drivers 550
• CUDA/cuDNN version (if applicable): CUDA 12.1
• 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): I'll add it later.
• Any other relevant information:

Additional context

TVM Unity Hash Tag

USE_NVTX: OFF
USE_GTEST: AUTO
SUMMARIZE: OFF
TVM_DEBUG_WITH_ABI_CHANGE: OFF
USE_IOS_RPC: OFF
USE_MSC: OFF
USE_ETHOSU:
CUDA_VERSION: 12.1
USE_LIBBACKTRACE: AUTO
DLPACK_PATH: 3rdparty/dlpack/include
USE_TENSORRT_CODEGEN: OFF
USE_THRUST: ON
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
TVM_LOG_BEFORE_THROW: OFF
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_MSCCL: OFF
USE_VITIS_AI: OFF
USE_MLIR: OFF
USE_RCCL: OFF
USE_LLVM: llvm-config --ignore-libllvm --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: ON
USE_ROCBLAS: OFF
GIT_COMMIT_HASH: f5f048bbd71513f087799f987019e3931f68a6d9
USE_VULKAN: ON
USE_RUST_EXT: OFF
USE_CUTLASS: ON
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: 2024-06-03 09:18:53 -0400
USE_HEXAGON_SDK: /path/to/sdk
USE_BLAS: none
USE_ETHOSN: OFF
USE_LIBTORCH: OFF
USE_RANDOM: ON
USE_CUDA: ON
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_MRVL: OFF
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_FLASHINFER: ON
USE_CUBLAS: ON
USE_METAL: OFF
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: /opt/rh/gcc-toolset-11/root/usr/bin/c++
HIDE_PRIVATE_SYMBOLS: ON

[MasterJH5574]

Thank you for reporting this! As I am trying to see what happens under the hood, I am wondering if you are able to run other models other than Qwen 72B with multi GPUs?

[alphaarea]

Thank you for reporting this! As I am trying to see what happens under the hood, I am wondering if you are able to run other models other than Qwen 72B with multi GPUs?

I can ran llama3-70b and qwen1.5-72b smoothly with the same setup.

A similar error is encountered when running Phi-3-mini-4k-instruct-q0f32, where the model terminates unexpectedly about 2/3 of the way through loading. The following error message is displayed here

[13:29:00] /workspace/tvm/src/runtime/memory/pooled_allocator.h:65: Warning: PooledAllocator got InternalError during allocation: InternalError: Check failed: (e == cudaSuccess || e == cudaErrorCudartUnloading) is false: CUDA: an illegal memory access was encountered
[13:29:00] /workspace/tvm/src/runtime/memory/pooled_allocator.h:66: Warning: Trying to release all unused memory and reallocate...
Aborted (core dumped)

[MasterJH5574]

@alphaarea Thanks for sharing the issue of CUDA: an illegal memory access was encountered which is really helpful here.

To try to get more insight, could you do me a favor? compute-sanitizer is a tool provided by CUDA that can check and report illegal memory accesses. Could you run the following command for both Phi-3-mini-4k-instruct-q0f32 and the Qwen2-72B?

compute-sanitizer --tool memcheck --print-limit=5 --target-processes all --launch-timeout 7200 mlc_llm chat --model-lib $MODEL_OUTPUT/$MODEL_LIB $MODEL_OUTPUT

It may take a bit longer time than directly running mlc_llm chat due to the memory access check. The expectation is that for the Phi-3 model, compute-sanitizer should report a stack trace showing where the illegal memory accesses happen. For the Qwen2 model, given we are not yet confirmed whether it fails due to illegal memory access on your side, it may or may not report the stack trace.

[alphaarea]

@MasterJH5574 Thank you for your reply.

I ran memcheck as you instructed, and Phi-3-mini-4k-instruct outputs a more detailed error message. But on Qwen2-72B-Instruct nothing is output, and I've been waiting for an hour with no change.

The detailed output is as follows

Phi-3-mini-4k-instruct

(mlc-llm-20240607) alphaarea@kcores-p100:~/models$ compute-sanitizer --tool memcheck --print-limit=5 --target-processes all --launch-timeout 7200 mlc_llm chat --model-lib ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so ./Phi-3-mini-4k-instruct-q0f32/
========= COMPUTE-SANITIZER
[2024-06-10 03:20:08] INFO auto_device.py:79: Found device: cuda:0
[2024-06-10 03:20:08] INFO auto_device.py:79: Found device: cuda:1
[2024-06-10 03:20:08] INFO auto_device.py:79: Found device: cuda:2
[2024-06-10 03:20:08] INFO auto_device.py:79: Found device: cuda:3
[2024-06-10 03:20:08] INFO auto_device.py:79: Found device: cuda:4
[2024-06-10 03:20:08] INFO auto_device.py:79: Found device: cuda:5
[2024-06-10 03:20:08] INFO auto_device.py:79: Found device: cuda:6
[2024-06-10 03:20:08] INFO auto_device.py:79: Found device: cuda:7
[2024-06-10 03:20:09] INFO auto_device.py:88: Not found device: rocm:0
[2024-06-10 03:20:10] INFO auto_device.py:88: Not found device: metal:0
[2024-06-10 03:20:13] INFO auto_device.py:79: Found device: vulkan:0
[2024-06-10 03:20:13] INFO auto_device.py:79: Found device: vulkan:1
[2024-06-10 03:20:13] INFO auto_device.py:79: Found device: vulkan:2
[2024-06-10 03:20:13] INFO auto_device.py:79: Found device: vulkan:3
[2024-06-10 03:20:13] INFO auto_device.py:79: Found device: vulkan:4
[2024-06-10 03:20:13] INFO auto_device.py:79: Found device: vulkan:5
[2024-06-10 03:20:13] INFO auto_device.py:79: Found device: vulkan:6
[2024-06-10 03:20:13] INFO auto_device.py:79: Found device: vulkan:7
[2024-06-10 03:20:14] INFO auto_device.py:88: Not found device: opencl:0
[2024-06-10 03:20:14] INFO auto_device.py:35: Using device: cuda:0
[2024-06-10 03:20:14] INFO engine_base.py:143: Using library model: ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so
[03:20:54] /workspace/mlc-llm/cpp/serve/config.cc:649: Under mode "local", max batch size will be set to 1, max KV cache token capacity will be set to 4096, prefill chunk size will be set to 2048.
[03:20:54] /workspace/mlc-llm/cpp/serve/config.cc:649: Under mode "interactive", max batch size will be set to 1, max KV cache token capacity will be set to 4096, prefill chunk size will be set to 2048.
[03:20:54] /workspace/mlc-llm/cpp/serve/config.cc:649: Under mode "server", max batch size will be set to 1, max KV cache token capacity will be set to 4096, prefill chunk size will be set to 2048.
[03:20:54] /workspace/mlc-llm/cpp/serve/config.cc:729: The actual engine mode is "interactive". So max batch size is 1, max KV cache token capacity is 4096, prefill chunk size is 2048.
[03:20:54] /workspace/mlc-llm/cpp/serve/config.cc:734: Estimated total single GPU memory usage: 2873.870 MB (Parameters: 2480.262 MB. KVCache: 246.083 MB. Temporary buffer: 147.525 MB). The actual usage might be slightly larger than the estimated number.
[03:20:54] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #0] Loading model to device: cuda:0
[03:20:54] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #1] Loading model to device: cuda:1
[03:20:54] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #2] Loading model to device: cuda:2
[03:20:54] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #3] Loading model to device: cuda:3
[03:20:54] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #7] Loading model to device: cuda:7
[03:20:54] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #6] Loading model to device: cuda:6
[03:20:54] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:165: Loading parameters...
[03:20:54] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #4] Loading model to device: cuda:4
[03:20:54] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #5] Loading model to device: cuda:5
[==================================================================================================>]  [195/195]
[03:21:45] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:184: Loading done. Time used: Loading 20.370 s Preprocessing 30.110 s.
You can use the following special commands:
  /help               print the special commands
  /exit               quit the cli
  /stats              print out stats of last request (token/sec)
  /metrics            print out full engine metrics
  /reset              restart a fresh chat
  /set [overrides]    override settings in the generation config. For example,
                      `/set temperature=0.5;top_p=0.8;seed=23;max_tokens=100;stop=str1,str2`
                      Note: Separate stop words in the `stop` option with commas (,).
  Multi-line input: Use escape+enter to start a new line.

========= Invalid __shared__ write of size 8 bytes
=========     at 0xa978 in batch_prefill_ragged_kv_kernel
=========     by thread (8,4,0) in block (0,0,0)
=========     Address 0x4cc0 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x2ef3d0]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:tvm::runtime::CUDAWrappedFunc::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*, void**) const [clone .isra.0] [0x343371f]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::detail::PackFuncVoidAddr_<0, tvm::runtime::CUDAWrappedFunc>(tvm::runtime::CUDAWrappedFunc, std::vector<tvm::runtime::detail::ArgConvertCode, std::allocator<tvm::runtime::detail::ArgConvertCode> > const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x3435e15]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:TVMFuncCall [0x32a636a]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:./IRModule.CodeGenLLVM:0:batch_prefill_ragged_kv_compute_ [0x1fc48]
=========                in ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so
=========     Host Frame:./IRModule.CodeGenLLVM:0:batch_prefill_ragged_kv [0xb285]
=========                in ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::WrapPackedFunc(int (*)(TVMValue*, int*, int, TVMValue*, int*, void*), tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x32eddb5]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::PagedAttentionKVCacheObj::AttentionInternal(long, tvm::runtime::NDArray, tvm::runtime::NDArray, tvm::runtime::NDArray, tvm::runtime::NDArray, double) [0x337bba4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::PagedAttentionKVCacheObj::AttentionWithFusedQKV(long, tvm::runtime::NDArray, tvm::runtime::Optional<tvm::runtime::NDArray>, tvm::runtime::NDArray, double) [0x337ced8]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::TypedPackedFunc<void (tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)>::AssignTypedLambda<tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}>(tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMArgs const&) const [clone .constprop.0] [0x3343f8d]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<void (tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)>::AssignTypedLambda<tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}>(tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, std::allocator<char>, tvm::runtime::TVMArgs const&) [0x3344052]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::RunInstrCall(tvm::runtime::relax_vm::VMFrame*, tvm::runtime::relax_vm::Instruction) [0x3388f2a]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::RunLoop() [0x33875a4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::InvokeBytecode(long, std::vector<tvm::runtime::TVMRetValue, std::allocator<tvm::runtime::TVMRetValue> > const&) [0x3387bf5]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::relax_vm::VirtualMachineImpl::GetClosureInternal(tvm::runtime::String const&, bool)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x33880ed]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::InvokeClosurePacked(tvm::runtime::ObjectRef const&, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x3385b69]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::DiscoWorker::Impl::CallPacked(tvm::runtime::DiscoWorker*, long, tvm::runtime::PackedFunc, tvm::runtime::TVMArgs const&) [0x32c0ee8]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::DiscoWorker::Impl::MainLoop(tvm::runtime::DiscoWorker*) [0x32c1554]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:/opt/conda/conda-bld/gcc-compiler_1654084175708/work/gcc/libstdc++-v3/src/c++11/thread.cc:84:execute_native_thread_routine [0xdbbf4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/bin/../lib/libstdc++.so.6
=========     Host Frame: [0x94ac3]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame: [0x126850]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========
========= Invalid __shared__ write of size 8 bytes
=========     at 0xa978 in batch_prefill_ragged_kv_kernel
=========     by thread (9,4,0) in block (0,0,0)
=========     Address 0x4cc8 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x2ef3d0]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:tvm::runtime::CUDAWrappedFunc::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*, void**) const [clone .isra.0] [0x343371f]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::detail::PackFuncVoidAddr_<0, tvm::runtime::CUDAWrappedFunc>(tvm::runtime::CUDAWrappedFunc, std::vector<tvm::runtime::detail::ArgConvertCode, std::allocator<tvm::runtime::detail::ArgConvertCode> > const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x3435e15]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:TVMFuncCall [0x32a636a]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:./IRModule.CodeGenLLVM:0:batch_prefill_ragged_kv_compute_ [0x1fc48]
=========                in ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so
=========     Host Frame:./IRModule.CodeGenLLVM:0:batch_prefill_ragged_kv [0xb285]
=========                in ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::WrapPackedFunc(int (*)(TVMValue*, int*, int, TVMValue*, int*, void*), tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x32eddb5]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::PagedAttentionKVCacheObj::AttentionInternal(long, tvm::runtime::NDArray, tvm::runtime::NDArray, tvm::runtime::NDArray, tvm::runtime::NDArray, double) [0x337bba4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::PagedAttentionKVCacheObj::AttentionWithFusedQKV(long, tvm::runtime::NDArray, tvm::runtime::Optional<tvm::runtime::NDArray>, tvm::runtime::NDArray, double) [0x337ced8]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::TypedPackedFunc<void (tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)>::AssignTypedLambda<tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}>(tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMArgs const&) const [clone .constprop.0] [0x3343f8d]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<void (tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)>::AssignTypedLambda<tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}>(tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, std::allocator<char>, tvm::runtime::TVMArgs const&) [0x3344052]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::RunInstrCall(tvm::runtime::relax_vm::VMFrame*, tvm::runtime::relax_vm::Instruction) [0x3388f2a]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::RunLoop() [0x33875a4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::InvokeBytecode(long, std::vector<tvm::runtime::TVMRetValue, std::allocator<tvm::runtime::TVMRetValue> > const&) [0x3387bf5]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::relax_vm::VirtualMachineImpl::GetClosureInternal(tvm::runtime::String const&, bool)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x33880ed]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::InvokeClosurePacked(tvm::runtime::ObjectRef const&, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x3385b69]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::DiscoWorker::Impl::CallPacked(tvm::runtime::DiscoWorker*, long, tvm::runtime::PackedFunc, tvm::runtime::TVMArgs const&) [0x32c0ee8]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::DiscoWorker::Impl::MainLoop(tvm::runtime::DiscoWorker*) [0x32c1554]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:/opt/conda/conda-bld/gcc-compiler_1654084175708/work/gcc/libstdc++-v3/src/c++11/thread.cc:84:execute_native_thread_routine [0xdbbf4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/bin/../lib/libstdc++.so.6
=========     Host Frame: [0x94ac3]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame: [0x126850]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========
========= Invalid __shared__ write of size 8 bytes
=========     at 0xa978 in batch_prefill_ragged_kv_kernel
=========     by thread (10,4,0) in block (0,0,0)
=========     Address 0x4cd0 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x2ef3d0]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:tvm::runtime::CUDAWrappedFunc::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*, void**) const [clone .isra.0] [0x343371f]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::detail::PackFuncVoidAddr_<0, tvm::runtime::CUDAWrappedFunc>(tvm::runtime::CUDAWrappedFunc, std::vector<tvm::runtime::detail::ArgConvertCode, std::allocator<tvm::runtime::detail::ArgConvertCode> > const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x3435e15]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:TVMFuncCall [0x32a636a]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:./IRModule.CodeGenLLVM:0:batch_prefill_ragged_kv_compute_ [0x1fc48]
=========                in ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so
=========     Host Frame:./IRModule.CodeGenLLVM:0:batch_prefill_ragged_kv [0xb285]
=========                in ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::WrapPackedFunc(int (*)(TVMValue*, int*, int, TVMValue*, int*, void*), tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x32eddb5]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::PagedAttentionKVCacheObj::AttentionInternal(long, tvm::runtime::NDArray, tvm::runtime::NDArray, tvm::runtime::NDArray, tvm::runtime::NDArray, double) [0x337bba4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::PagedAttentionKVCacheObj::AttentionWithFusedQKV(long, tvm::runtime::NDArray, tvm::runtime::Optional<tvm::runtime::NDArray>, tvm::runtime::NDArray, double) [0x337ced8]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::TypedPackedFunc<void (tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)>::AssignTypedLambda<tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}>(tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMArgs const&) const [clone .constprop.0] [0x3343f8d]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<void (tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)>::AssignTypedLambda<tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}>(tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, std::allocator<char>, tvm::runtime::TVMArgs const&) [0x3344052]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::RunInstrCall(tvm::runtime::relax_vm::VMFrame*, tvm::runtime::relax_vm::Instruction) [0x3388f2a]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::RunLoop() [0x33875a4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::InvokeBytecode(long, std::vector<tvm::runtime::TVMRetValue, std::allocator<tvm::runtime::TVMRetValue> > const&) [0x3387bf5]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::relax_vm::VirtualMachineImpl::GetClosureInternal(tvm::runtime::String const&, bool)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x33880ed]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::InvokeClosurePacked(tvm::runtime::ObjectRef const&, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x3385b69]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::DiscoWorker::Impl::CallPacked(tvm::runtime::DiscoWorker*, long, tvm::runtime::PackedFunc, tvm::runtime::TVMArgs const&) [0x32c0ee8]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::DiscoWorker::Impl::MainLoop(tvm::runtime::DiscoWorker*) [0x32c1554]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:/opt/conda/conda-bld/gcc-compiler_1654084175708/work/gcc/libstdc++-v3/src/c++11/thread.cc:84:execute_native_thread_routine [0xdbbf4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/bin/../lib/libstdc++.so.6
=========     Host Frame: [0x94ac3]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame: [0x126850]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========
========= Invalid __shared__ write of size 8 bytes
=========     at 0xa978 in batch_prefill_ragged_kv_kernel
=========     by thread (11,4,0) in block (0,0,0)
=========     Address 0x4cd8 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x2ef3d0]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:tvm::runtime::CUDAWrappedFunc::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*, void**) const [clone .isra.0] [0x343371f]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::detail::PackFuncVoidAddr_<0, tvm::runtime::CUDAWrappedFunc>(tvm::runtime::CUDAWrappedFunc, std::vector<tvm::runtime::detail::ArgConvertCode, std::allocator<tvm::runtime::detail::ArgConvertCode> > const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x3435e15]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:TVMFuncCall [0x32a636a]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:./IRModule.CodeGenLLVM:0:batch_prefill_ragged_kv_compute_ [0x1fc48]
=========                in ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so
=========     Host Frame:./IRModule.CodeGenLLVM:0:batch_prefill_ragged_kv [0xb285]
=========                in ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::WrapPackedFunc(int (*)(TVMValue*, int*, int, TVMValue*, int*, void*), tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x32eddb5]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::PagedAttentionKVCacheObj::AttentionInternal(long, tvm::runtime::NDArray, tvm::runtime::NDArray, tvm::runtime::NDArray, tvm::runtime::NDArray, double) [0x337bba4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::PagedAttentionKVCacheObj::AttentionWithFusedQKV(long, tvm::runtime::NDArray, tvm::runtime::Optional<tvm::runtime::NDArray>, tvm::runtime::NDArray, double) [0x337ced8]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::TypedPackedFunc<void (tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)>::AssignTypedLambda<tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}>(tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMArgs const&) const [clone .constprop.0] [0x3343f8d]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<void (tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)>::AssignTypedLambda<tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}>(tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, std::allocator<char>, tvm::runtime::TVMArgs const&) [0x3344052]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::RunInstrCall(tvm::runtime::relax_vm::VMFrame*, tvm::runtime::relax_vm::Instruction) [0x3388f2a]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::RunLoop() [0x33875a4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::InvokeBytecode(long, std::vector<tvm::runtime::TVMRetValue, std::allocator<tvm::runtime::TVMRetValue> > const&) [0x3387bf5]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::relax_vm::VirtualMachineImpl::GetClosureInternal(tvm::runtime::String const&, bool)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x33880ed]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::InvokeClosurePacked(tvm::runtime::ObjectRef const&, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x3385b69]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::DiscoWorker::Impl::CallPacked(tvm::runtime::DiscoWorker*, long, tvm::runtime::PackedFunc, tvm::runtime::TVMArgs const&) [0x32c0ee8]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::DiscoWorker::Impl::MainLoop(tvm::runtime::DiscoWorker*) [0x32c1554]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:/opt/conda/conda-bld/gcc-compiler_1654084175708/work/gcc/libstdc++-v3/src/c++11/thread.cc:84:execute_native_thread_routine [0xdbbf4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/bin/../lib/libstdc++.so.6
=========     Host Frame: [0x94ac3]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame: [0x126850]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========
========= Invalid __shared__ write of size 8 bytes
=========     at 0xa978 in batch_prefill_ragged_kv_kernel
=========     by thread (12,4,0) in block (0,0,0)
=========     Address 0x4ce0 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x2ef3d0]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:tvm::runtime::CUDAWrappedFunc::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*, void**) const [clone .isra.0] [0x343371f]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::detail::PackFuncVoidAddr_<0, tvm::runtime::CUDAWrappedFunc>(tvm::runtime::CUDAWrappedFunc, std::vector<tvm::runtime::detail::ArgConvertCode, std::allocator<tvm::runtime::detail::ArgConvertCode> > const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x3435e15]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:TVMFuncCall [0x32a636a]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:./IRModule.CodeGenLLVM:0:batch_prefill_ragged_kv_compute_ [0x1fc48]
=========                in ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so
=========     Host Frame:./IRModule.CodeGenLLVM:0:batch_prefill_ragged_kv [0xb285]
=========                in ./Phi-3-mini-4k-instruct-q0f32/Phi-3-mini-4k-instruct-q0f32-cuda.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::WrapPackedFunc(int (*)(TVMValue*, int*, int, TVMValue*, int*, void*), tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x32eddb5]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::PagedAttentionKVCacheObj::AttentionInternal(long, tvm::runtime::NDArray, tvm::runtime::NDArray, tvm::runtime::NDArray, tvm::runtime::NDArray, double) [0x337bba4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::PagedAttentionKVCacheObj::AttentionWithFusedQKV(long, tvm::runtime::NDArray, tvm::runtime::Optional<tvm::runtime::NDArray>, tvm::runtime::NDArray, double) [0x337ced8]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::TypedPackedFunc<void (tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)>::AssignTypedLambda<tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}>(tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMArgs const&) const [clone .constprop.0] [0x3343f8d]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<void (tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)>::AssignTypedLambda<tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}>(tvm::runtime::relax_vm::__mk_TVM17::{lambda(tvm::runtime::relax_vm::AttentionKVCache, long, double, tvm::runtime::NDArray, tvm::runtime::NDArray)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, std::allocator<char>, tvm::runtime::TVMArgs const&) [0x3344052]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::RunInstrCall(tvm::runtime::relax_vm::VMFrame*, tvm::runtime::relax_vm::Instruction) [0x3388f2a]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::RunLoop() [0x33875a4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::InvokeBytecode(long, std::vector<tvm::runtime::TVMRetValue, std::allocator<tvm::runtime::TVMRetValue> > const&) [0x3387bf5]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::relax_vm::VirtualMachineImpl::GetClosureInternal(tvm::runtime::String const&, bool)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x33880ed]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::relax_vm::VirtualMachineImpl::InvokeClosurePacked(tvm::runtime::ObjectRef const&, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) [0x3385b69]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::DiscoWorker::Impl::CallPacked(tvm::runtime::DiscoWorker*, long, tvm::runtime::PackedFunc, tvm::runtime::TVMArgs const&) [0x32c0ee8]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:tvm::runtime::DiscoWorker::Impl::MainLoop(tvm::runtime::DiscoWorker*) [0x32c1554]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/lib/python3.11/site-packages/tvm/libtvm.so
=========     Host Frame:/opt/conda/conda-bld/gcc-compiler_1654084175708/work/gcc/libstdc++-v3/src/c++11/thread.cc:84:execute_native_thread_routine [0xdbbf4]
=========                in /home/alphaarea/.conda/envs/mlc-llm-20240607/bin/../lib/libstdc++.so.6
=========     Host Frame: [0x94ac3]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame: [0x126850]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========
========= Error: process didn't terminate successfully
========= Target application returned an error
========= ERROR SUMMARY: 1557 errors
========= ERROR SUMMARY: 1552 errors were not printed. Use --print-limit option to adjust the number of printed errors

Qwen2-72B-Instruct

(mlc-llm-20240607) alphaarea@kcores-p100:~/models$ compute-sanitizer --tool memcheck --print-limit=5 --target-processes all --launch-timeout 7200 mlc_llm chat --model-lib Qwen2-72B-Instruct-q4f16_1/Qwen2-72B-Instruct-q4f16_1-cuda.so ./Qwen2-72B-Instruct-q4f16_1/
========= COMPUTE-SANITIZER
[2024-06-10 04:19:54] INFO auto_device.py:79: Found device: cuda:0
[2024-06-10 04:19:54] INFO auto_device.py:79: Found device: cuda:1
[2024-06-10 04:19:54] INFO auto_device.py:79: Found device: cuda:2
[2024-06-10 04:19:54] INFO auto_device.py:79: Found device: cuda:3
[2024-06-10 04:19:54] INFO auto_device.py:79: Found device: cuda:4
[2024-06-10 04:19:54] INFO auto_device.py:79: Found device: cuda:5
[2024-06-10 04:19:54] INFO auto_device.py:79: Found device: cuda:6
[2024-06-10 04:19:54] INFO auto_device.py:79: Found device: cuda:7
[2024-06-10 04:19:55] INFO auto_device.py:88: Not found device: rocm:0
[2024-06-10 04:19:56] INFO auto_device.py:88: Not found device: metal:0
[2024-06-10 04:19:59] INFO auto_device.py:79: Found device: vulkan:0
[2024-06-10 04:19:59] INFO auto_device.py:79: Found device: vulkan:1
[2024-06-10 04:19:59] INFO auto_device.py:79: Found device: vulkan:2
[2024-06-10 04:19:59] INFO auto_device.py:79: Found device: vulkan:3
[2024-06-10 04:19:59] INFO auto_device.py:79: Found device: vulkan:4
[2024-06-10 04:19:59] INFO auto_device.py:79: Found device: vulkan:5
[2024-06-10 04:19:59] INFO auto_device.py:79: Found device: vulkan:6
[2024-06-10 04:19:59] INFO auto_device.py:79: Found device: vulkan:7
[2024-06-10 04:20:00] INFO auto_device.py:88: Not found device: opencl:0
[2024-06-10 04:20:00] INFO auto_device.py:35: Using device: cuda:0
[2024-06-10 04:20:00] INFO engine_base.py:143: Using library model: Qwen2-72B-Instruct-q4f16_1/Qwen2-72B-Instruct-q4f16_1-cuda.so
[04:20:40] /workspace/mlc-llm/cpp/serve/config.cc:649: Under mode "local", max batch size will be set to 1, max KV cache token capacity will be set to 8192, prefill chunk size will be set to 2048.
[04:20:40] /workspace/mlc-llm/cpp/serve/config.cc:649: Under mode "interactive", max batch size will be set to 1, max KV cache token capacity will be set to 32768, prefill chunk size will be set to 2048.
[04:20:40] /workspace/mlc-llm/cpp/serve/config.cc:649: Under mode "server", max batch size will be set to 1, max KV cache token capacity will be set to 32768, prefill chunk size will be set to 2048.
[04:20:40] /workspace/mlc-llm/cpp/serve/config.cc:729: The actual engine mode is "interactive". So max batch size is 1, max KV cache token capacity is 32768, prefill chunk size is 2048.
[04:20:40] /workspace/mlc-llm/cpp/serve/config.cc:734: Estimated total single GPU memory usage: 10032.489 MB (Parameters: 6052.961 MB. KVCache: 1337.180 MB. Temporary buffer: 2642.348 MB). The actual usage might be slightly larger than the estimated number.
[04:20:40] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #0] Loading model to device: cuda:0
[04:20:40] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #1] Loading model to device: cuda:1
[04:20:40] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #2] Loading model to device: cuda:2
[04:20:40] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #3] Loading model to device: cuda:3
[04:20:40] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #7] Loading model to device: cuda:7
[04:20:40] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #4] Loading model to device: cuda:4[04:20:40] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #5] Loading model to device: cuda:5

[04:20:40] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:140: [Worker #6] Loading model to device: cuda:6
[04:20:40] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:165: Loading parameters...
[>                                                                                                  ]  [2/885

[MasterJH5574]

@alphaarea Thank you for the valuable information. So they are two independent issues, and we'll look into those.

I just reproduced the error of Qwen2-72B on 8 V100s. Here is the stack trace for your information. It's interesting to me that your side does not show any info. I'll report back with updates.

[11:19:45] /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:165: Loading parameters...
terminate called after throwing an instance of 'tvm::runtime::Error'                                ]  [2/885]
  what():  TVMError: Assert fail: T.Cast("int32", _shard_mlp_down_q_weight_var_w_shape[1]) == 3712, Argument _shard_mlp_down_q_weight.var_w.shape[1] has an unsatisfied constraint: 3712 == T.Cast("int32", _shard_mlp_down_q_weight_var_w_shape[1])
Stack trace:
  0: TVMThrowLastError.cold
  1: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::WrapPackedFunc(int (*)(TVMValue*, int*, int, TVMValue*, int*, void*), tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)
  2: mlc::llm::loader::PreprocessorPool::Apply(tvm::runtime::NDArray, mlc::llm::ModelMetadata::Param const&) const
        at /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:76
  3: mlc::llm::loader::BroadcastOrShardAndScatter(tvm::runtime::NDArray, mlc::llm::ModelMetadata::Param const&, int, mlc::llm::loader::PreprocessorPool const&)
        at /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:105
  4: mlc::llm::loader::LoadMultiGPU(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, tvm::runtime::Module, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&)
        at /workspace/mlc-llm/cpp/loader/multi_gpu_loader.cc:179
  5: tvm::runtime::DiscoWorker::Impl::CallPacked(tvm::runtime::DiscoWorker*, long, tvm::runtime::PackedFunc, tvm::runtime::TVMArgs const&)
  6: tvm::runtime::DiscoWorker::Impl::MainLoop(tvm::runtime::DiscoWorker*)
  7: execute_native_thread_routine
        at ../../../../../libstdc++-v3/src/c++11/thread.cc:104
  8: start_thread
  9: __GI___clone
  10: 0xffffffffffffffff

[MasterJH5574]

Hi @alphaarea, just want to share a quick update here. We identified the failure reason of Qwen2-72B on 8 GPUs for quantization q4f16_1.

The q4f16_1 in MLC LLM is group quantization with group size 32, which means we group 32 consecutive fp16 numbers together to compute the group scale. The MLP of Qwen2-72B model has intermediate size 29568, which has 29568 // 32 = 924 groups under q4f16_1. However, 924 is not a multiple of 8 (924 / 8 = 115.5), which means we cannot evenly distribute the 924 groups across 8 GPUs. So the combination of “Qwen2-72B” + “q4f16_1” + “8 GPUs” does not work well here.

On one hand, we will improve the error reporting to have clear error message when running mlc_llm compile to notify people about such inability. On the other hand, if you still want to run quantized Qwen2-72B on 8 GPUs, at this moment what you can do to work around is to find the file quantization/quantization.py on your local side, change the group size from 32 to 16, and run mlc_llm convert_weight and mlc_llm compile again. Since under group size 16 there will be 29568 // 16 = 1848 groups which is a multiple of 8, it's expected to work for 8 GPUs.

mlc-llm/python/mlc_llm/quantization/quantization.py

Line 77 in 4234262

group_size=32,

We'll further look into the Phi-3 issue you shared here and report the updates.

[MasterJH5574]

Fixed error reporting in #2566.