[Bug]: Crash with --enable-prefix-caching enabled

[eByteTheDust]

Your current environment

Collecting environment information...
PyTorch version: 2.1.2+cu121
Is debug build: False
CUDA used to build PyTorch: 12.1
ROCM used to build PyTorch: N/A

OS: Ubuntu 22.04.4 LTS (x86_64)
GCC version: (Ubuntu 11.4.0-1ubuntu1~22.04) 11.4.0
Clang version: Could not collect
CMake version: version 3.29.0
Libc version: glibc-2.35

Python version: 3.11.7 (main, Dec 15 2023, 18:12:31) [GCC 11.2.0] (64-bit runtime)
Python platform: Linux-6.5.0-1017-azure-x86_64-with-glibc2.35
Is CUDA available: True
CUDA runtime version: Could not collect
CUDA_MODULE_LOADING set to: LAZY
GPU models and configuration: GPU 0: Tesla T4
Nvidia driver version: 535.161.07
cuDNN version: Could not collect
HIP runtime version: N/A
MIOpen runtime version: N/A
Is XNNPACK available: True

CPU:
Architecture:                       x86_64
CPU op-mode(s):                     32-bit, 64-bit
Address sizes:                      48 bits physical, 48 bits virtual
Byte Order:                         Little Endian
CPU(s):                             4
On-line CPU(s) list:                0-3
Vendor ID:                          AuthenticAMD
Model name:                         AMD EPYC 7V12 64-Core Processor
CPU family:                         23
Model:                              49
Thread(s) per core:                 1
Core(s) per socket:                 4
Socket(s):                          1
Stepping:                           0
BogoMIPS:                           4890.88
Flags:                              fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 ht syscall nx mmxext fxsr_opt pdpe1gb rdtscp lm constant_tsc rep_good nopl tsc_reliable nonstop_tsc cpuid extd_apicid aperfmperf pni pclmulqdq ssse3 fma cx16 sse4_1 sse4_2 movbe popcnt aes xsave avx f16c rdrand hypervisor lahf_lm cmp_legacy cr8_legacy abm sse4a misalignsse 3dnowprefetch osvw topoext perfctr_core ssbd vmmcall fsgsbase bmi1 avx2 smep bmi2 rdseed adx smap clflushopt clwb sha_ni xsaveopt xsavec xgetbv1 clzero xsaveerptr rdpru arat umip rdpid
Hypervisor vendor:                  Microsoft
Virtualization type:                full
L1d cache:                          128 KiB (4 instances)
L1i cache:                          128 KiB (4 instances)
L2 cache:                           2 MiB (4 instances)
L3 cache:                           16 MiB (1 instance)
NUMA node(s):                       1
NUMA node0 CPU(s):                  0-3
Vulnerability Gather data sampling: Not affected
Vulnerability Itlb multihit:        Not affected
Vulnerability L1tf:                 Not affected
Vulnerability Mds:                  Not affected
Vulnerability Meltdown:             Not affected
Vulnerability Mmio stale data:      Not affected
Vulnerability Retbleed:             Mitigation; untrained return thunk; SMT disabled
Vulnerability Spec rstack overflow: Vulnerable: Safe RET, no microcode
Vulnerability Spec store bypass:    Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1:           Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2:           Mitigation; Retpolines, STIBP disabled, RSB filling, PBRSB-eIBRS Not affected
Vulnerability Srbds:                Not affected
Vulnerability Tsx async abort:      Not affected

Versions of relevant libraries:
[pip3] mypy-extensions==1.0.0
[pip3] numpy==1.26.4
[pip3] torch==2.1.2
[pip3] torchaudio==2.1.2+cu121
[pip3] torchvision==0.16.2+cu121
[pip3] triton==2.1.0
[conda] blas                      1.0                         mkl
[conda] ffmpeg                    4.3                  hf484d3e_0    pytorch
[conda] libjpeg-turbo             2.0.0                h9bf148f_0    pytorch
[conda] mkl                       2023.1.0         h213fc3f_46344
[conda] mkl-service               2.4.0           py311h5eee18b_1
[conda] mkl_fft                   1.3.8           py311h5eee18b_0
[conda] mkl_random                1.2.4           py311hdb19cb5_0
[conda] numpy                     1.26.4                   pypi_0    pypi
[conda] pytorch-cuda              12.1                 ha16c6d3_5    pytorch
[conda] pytorch-mutex             1.0                        cuda    pytorch
[conda] torch                     2.1.2                    pypi_0    pypi
[conda] torchaudio                2.1.2+cu121              pypi_0    pypi
[conda] torchvision               0.16.2+cu121             pypi_0    pypi
[conda] triton                    2.1.0                    pypi_0    pypiROCM Version: Could not collect
Neuron SDK Version: N/A
vLLM Version: 0.4.0.post1
vLLM Build Flags:
CUDA Archs: Not Set; ROCm: Disabled; Neuron: Disabled
GPU Topology:
GPU0    NIC0    CPU Affinity    NUMA Affinity   GPU NUMA ID
GPU0     X      SYS     0-3     0               N/A
NIC0    SYS      X

Legend:

  X    = Self
  SYS  = Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI)
  NODE = Connection traversing PCIe as well as the interconnect between PCIe Host Bridges within a NUMA node
  PHB  = Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU)
  PXB  = Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge)
  PIX  = Connection traversing at most a single PCIe bridge
  NV#  = Connection traversing a bonded set of # NVLinks

NIC Legend:

  NIC0: mlx5_0

🐛 Describe the bug

when 'enable-prefix-caching' is on, VLLM always crash when the second message is sent to the server. The first message always works and it always crashes on the second (I guess when it is trying to use cache). I'm using ray and v0.4.0.post1. When I remove 'enable-prefix-caching' is works well.
python3 -m vllm.entrypoints.openai.api_server --model /maindir/Nous-Hermes-2-SOLAR-10.7B --tensor-parallel-size 2 --dtype half --gpu-memory-utilization 0.96 --max-model-len 4096 --enforce-eager --worker-use-ray --load-format auto --disable-log-stats --max-context-len-to-capture 4096 --enable-prefix-caching.

python3: /project/lib/Analysis/Allocation.cpp:40: std::pair<llvm::SmallVector, llvm::SmallVector > mlir::triton::getCvtOrder(mlir::Attribute, mlir::Attribute): Assertion `!(srcMmaLayout && dstMmaLayout) && "Unexpected mma -> mma layout conversion"' failed.
*** SIGABRT received at time=1712674906 on cpu 3 ***
PC: @ 0x7e0584c969fc (unknown) pthread_kill
@ 0x7e0584c42520 (unknown) (unknown)
[2024-04-09 11:01:46,817 E 19542 19712] logging.cc:361: *** SIGABRT received at time=1712674906 on cpu 3 ***
[2024-04-09 11:01:46,817 E 19542 19712] logging.cc:361: PC: @ 0x7e0584c969fc (unknown) pthread_kill
[2024-04-09 11:01:46,817 E 19542 19712] logging.cc:361: @ 0x7e0584c42520 (unknown) (unknown)
Fatal Python error: Aborted

Stack (most recent call first):
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/triton/compiler/compiler.py", line 107 in ttgir_to_llir
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/triton/compiler/compiler.py", line 385 in
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/triton/compiler/compiler.py", line 476 in compile
File "", line 63 in _fwd_kernel
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/vllm/attention/ops/prefix_prefill.py", line 699 in context_attention_fwd
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/utils/_contextlib.py", line 115 in decorate_context
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/vllm/attention/ops/paged_attn.py", line 178 in forward_prefix
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/vllm/attention/backends/xformers.py", line 262 in forward
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/vllm/attention/layer.py", line 46 in forward
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1527 in _call_impl
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1518 in _wrapped_call_impl
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/vllm/model_executor/models/llama.py", line 156 in forward
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1527 in _call_impl
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1518 in _wrapped_call_impl
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/vllm/model_executor/models/llama.py", line 213 in forward
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1527 in _call_impl
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1518 in _wrapped_call_impl
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/vllm/model_executor/models/llama.py", line 271 in forward
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1527 in _call_impl
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1518 in _wrapped_call_impl
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/vllm/model_executor/models/llama.py", line 345 in forward
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1527 in _call_impl
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1518 in _wrapped_call_impl
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/vllm/worker/model_runner.py", line 663 in execute_model
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/utils/_contextlib.py", line 115 in decorate_context
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/vllm/worker/worker.py", line 221 in execute_model
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/site-packages/torch/utils/_contextlib.py", line 115 in decorate_context
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/concurrent/futures/thread.py", line 58 in run
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/concurrent/futures/thread.py", line 83 in _worker
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/threading.py", line 982 in run
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/threading.py", line 1045 in _bootstrap_inner
File "/home/server11/miniconda3/envs/maindir/lib/python3.11/threading.py", line 1002 in _bootstrap

Extension modules: numpy.core._multiarray_umath, numpy.core._multiarray_tests, numpy.linalg._umath_linalg, numpy.fft._pocketfft_internal, numpy.random._common, numpy.random.bit_generator, numpy.random._bounded_integers, numpy.random._mt19937, numpy.random.mtrand, numpy.random._philox, numpy.random._pcg64, numpy.random._sfc64, numpy.random._generator, torch._C, torch._C._fft, torch._C._linalg, torch._C._nested, torch._C._nn, torch._C._sparse, torch._C._special, gmpy2.gmpy2, _brotli, zstandard.backend_c, charset_normalizer.md, yaml._yaml, sentencepiece._sentencepiece, psutil._psutil_linux, psutil._psutil_posix, msgpack._cmsgpack, google._upb._message, setproctitle, uvloop.loop, ray._raylet, regex._regex, scipy._lib._ccallback_c, numba.core.typeconv._typeconv, numba._helperlib, numba._dynfunc, numba._dispatcher, numba.core.runtime._nrt_python, numba.np.ufunc._internal, numba.experimental.jitclass._box, markupsafe._speedups, pyarrow.lib, pyarrow._hdfsio, pyarrow._json, PIL._imaging, scipy.linalg._fblas, scipy.linalg._flapack, scipy.linalg.cython_lapack, scipy.linalg._cythonized_array_utils, scipy.linalg._solve_toeplitz, scipy.linalg._flinalg, scipy.linalg._decomp_lu_cython, scipy.linalg._matfuncs_sqrtm_triu, scipy.linalg.cython_blas, scipy.linalg._matfuncs_expm, scipy.linalg._decomp_update, scipy.sparse._sparsetools, _csparsetools, scipy.sparse._csparsetools, scipy.sparse.linalg._dsolve._superlu, scipy.sparse.linalg._eigen.arpack._arpack, scipy.sparse.csgraph._tools, scipy.sparse.csgraph._shortest_path, scipy.sparse.csgraph._traversal, scipy.sparse.csgraph._min_spanning_tree, scipy.sparse.csgraph._flow, scipy.sparse.csgraph._matching, scipy.sparse.csgraph._reordering, scipy.optimize._minpack2, scipy.optimize._group_columns, scipy._lib.messagestream, scipy.optimize._trlib._trlib, scipy.optimize._lbfgsb, _moduleTNC, scipy.optimize._moduleTNC, scipy.optimize._cobyla, scipy.optimize._slsqp, scipy.optimize._minpack, scipy.optimize._lsq.givens_elimination, scipy.optimize._zeros, scipy.optimize._highs.cython.src._highs_wrapper, scipy.optimize._highs._highs_wrapper, scipy.optimize._highs.cython.src._highs_constants, scipy.optimize._highs._highs_constants, scipy.linalg._interpolative, scipy.optimize._bglu_dense, scipy.optimize._lsap, scipy.spatial._ckdtree, scipy.spatial._qhull, scipy.spatial._voronoi, scipy.spatial._distance_wrap, scipy.spatial._hausdorff, scipy.special._ufuncs_cxx, scipy.special._ufuncs, scipy.special._specfun, scipy.special._comb, scipy.special._ellip_harm_2, scipy.spatial.transform._rotation, scipy.optimize._direct, httptools.parser.parser, httptools.parser.url_parser, websockets.speedups (total: 104)
./solar-ray.sh: line 9: 19542 Aborted (core dumped) python3 -m vllm.entrypoints.openai.api_server --model /maindir/Nous-Hermes-2-SOLAR-10.7B --tensor-parallel-size 2 --dtype half --gpu-memory-utilization 0.96 --max-model-len 4096 --enforce-eager --worker-use-ray --load-format auto --disable-log-stats --max-context-len-to-capture 4096 --enable-prefix-caching

[robertgshaw2-redhat]

Right now, APC is not supported on T4 due to the attention kernels requiring > compute capability 8.0 (Ampere)

I added a check for this in #3903 with a better error message.

Please open a Feature Request for T4 support if you're interested or open a PR with T4 support if interested

[dichild]

upgrade the triton.
docker exec -it "name" /bin/bash
pip install triton --upgrade

[OUTHIM]

Same problem with Testla V100 32GB capability=7.0. Is it due to the low capability of my GPU?

[robertgshaw2-redhat]

Yes

[OUTHIM]

Thank you for your response @robertgshaw2-neuralmagic.

Does the limit only occur on the automatic-prefix-caching functionality? Can I surpass the limit of the capability using the prefix_pos argument in the previous versions?

[eByteTheDust]

I upgraded to v0.4.1 and prefix caching is working well on a T4. I mean I added the parameter and the server no longer crashes after the second message. It looks like it’s working

[robertgshaw2-redhat]

Oh Thanks @eByteTheDust

I think v0.4.1 had a triton upgrade which seems to have resolved the issues with the previous kernel for prefix attention

[OUTHIM]

I upgraded to v0.4.1 and prefix caching is working well on a T4. I mean I added the parameter and the server no longer crashes after the second message. It looks like it’s working

@eByteTheDust That sounds great! Can you provide the version of your triton for reference?
I am using vllm=0.4.1 and triton=2.2.0.
Maybe my triton version is out-dated!

[OUTHIM]

@eByteTheDust @robertgshaw2-neuralmagic Just an update:
Upgrading triton to 2.3.0 solves my problem.
It seems that it is not quite related to the GPU capability.
I share my configs here for you and others as a reference.

GPU: Tesla V100 capability=7.0.
triton=2.3.0
vllm=0.4.1

[biaochen]

I have the same issue, is there a solution ?
GPU: Tesla V100 capability=7.0
triton=2.3.0
vllm=0.4.2 or v0.5.0.post1 (neither works)

[shuaiyu5]

APC is not supported on T4 due to the attention kernels requiring > compute capability 8.0 (Ampere)

same issue，i've tried a range of vllm version with triton, none of them works. anyone to help, thanks!!!!

[shuaiyu5]

APC is not supported on T4 due to the attention kernels requiring > compute capability 8.0 (Ampere)

same issue，i've tried a range of vllm version with triton, none of them works. anyone to help, thanks!!!!

@robertgshaw2-neuralmagic

[robertgshaw2-redhat]

Can you share the specific request pattern you are sending. If I can reproduce the error I can try to fix it

[shuaiyu5]

Can you share the specific request pattern you are sending. If I can reproduce the error I can try to fix it

exactly same error as @eByteTheDust.
when 'enable-prefix-caching' is on, VLLM always crash when the second message is sent to the server. The first message always works and it always crashes on the second (I guess when it is trying to use cache). When I remove 'enable-prefix-caching' is works well.
the error information are
python3: /project/lib/Analysis/Allocation.cpp:40: std::pair<llvm::SmallVector, llvm::SmallVector > mlir::triton::getCvtOrder(mlir::Attribute, mlir::Attribute): Assertion `!(srcMmaLayout && dstMmaLayout) && "Unexpected mma -> mma layout conversion"' failed.

i've tried the method suggested as @OUTHIM did, but it didnt work for me.
i found the fatal information comes from triton, and i've tried 2.1.0、2.2.0、2.3.0 with proper vllm and torch, all not work.
is this the imcompatibility of prefix caching with v100??

thanks for your reply!!!! @robertgshaw2-neuralmagic

[OUTHIM]

@eByteTheDust @robertgshaw2-neuralmagic Just an update: Upgrading triton to 2.3.0 solves my problem. It seems that it is not quite related to the GPU capability. I share my configs here for you and others as a reference.

GPU: Tesla V100 capability=7.0. triton=2.3.0 vllm=0.4.1

An update from me:
Using this configuration, if the second input is a bit longer, it crashes. I guess there is still some incompatibility inside.
I give up using this functionality although it will be quite useful if it can work.

[mces89]

@robertgshaw2-neuralmagic @OUTHIM @shuaiyu5 are you able to solve this issue on v100? I'm using v100 with vllm 0.6.3.post1 pre-build docker, still get the same crash...

[OUTHIM]

@robertgshaw2-neuralmagic @OUTHIM @shuaiyu5 are you able to solve this issue on v100? I'm using v100 with vllm 0.6.3.post1 pre-build docker, still get the same crash...

No. It could run but when the prefix was a bit longer it crashed again.
I did not continue to explore this since then.