Your current environment
Environment info
Collecting environment information...
==============================
System Info
==============================
OS : Ubuntu 22.04.5 LTS (x86_64)
GCC version : (Ubuntu 11.4.0-1ubuntu1~22.04) 11.4.0
Clang version : Could not collect
CMake version : version 3.22.1
Libc version : glibc-2.35
==============================
PyTorch Info
==============================
PyTorch version : 2.13.0+cu130
Is debug build : False
CUDA used to build PyTorch : 13.0
ROCM used to build PyTorch : N/A
XPU used to build PyTorch : N/A
==============================
Python Environment
==============================
Python version : 3.12.3 | packaged by Anaconda, Inc. | (main, May 6 2024, 19:46:43) [GCC 11.2.0] (64-bit runtime)
Python platform : Linux-5.15.0-78-generic-x86_64-with-glibc2.35
==============================
CUDA / GPU Info
==============================
Is CUDA available : True
CUDA runtime version : 12.8.93
CUDA_MODULE_LOADING set to :
GPU models and configuration : GPU 0: NVIDIA RTX PRO 6000 Blackwell Server Edition
Nvidia driver version : 595.71.05
cuDNN version : Probably one of the following:
/usr/lib/x86_64-linux-gnu/libcudnn.so.9.8.0
/usr/lib/x86_64-linux-gnu/libcudnn_adv.so.9.8.0
/usr/lib/x86_64-linux-gnu/libcudnn_cnn.so.9.8.0
/usr/lib/x86_64-linux-gnu/libcudnn_engines_precompiled.so.9.8.0
/usr/lib/x86_64-linux-gnu/libcudnn_engines_runtime_compiled.so.9.8.0
/usr/lib/x86_64-linux-gnu/libcudnn_graph.so.9.8.0
/usr/lib/x86_64-linux-gnu/libcudnn_heuristic.so.9.8.0
/usr/lib/x86_64-linux-gnu/libcudnn_ops.so.9.8.0
HIP runtime version : N/A
MIOpen runtime version : N/A
Is XNNPACK available : False
==============================
vLLM Info
==============================
ROCM Version : Could not collect
vLLM Version : 0.27.1
vLLM Build Flags:
CUDA Archs: Not Set; ROCm: Disabled; XPU: Disabled
🐛 Describe the bug
Out-of-bounds attrIdxs in the C++ batch memcpy path (cache_kernels.cu)
Summary
The C++ batch-copy path in csrc/libtorch_stable/cache_kernels.cu passes a single size_t as the attrIdxs array to cuMemcpyBatchAsync / hipMemcpyBatchAsync, regardless of the descriptor count — the same contract violation fixed for the Python path in #53860.
The CUDA Driver API documentation for cuMemcpyBatchAsync specifies attrIdxs as an array of count elements (indices into attrs, each < numAttrs). Passing &attrs_idx where attrs_idx is one stack scalar while count is large (typically thousands of 64 KiB KV blocks per offload) leaves the driver reading past the end of that scalar — an out-of-bounds read / undefined behavior that currently happens to work only because adjacent stack bytes happen to read as tolerable values.
Location
Two sites, both in the batch-copy helper (around L130–L230 of csrc/libtorch_stable/cache_kernels.cu):
CUDA path (L169):
size_t attrs_idx = 0;
size_t fail_idx = 0;
...
CUresult result = batch_fn(..., static_cast<size_t>(cnt), &attr, &attrs_idx, 1,
&fail_idx, static_cast<CUstream>(stream));ROCm path (L196): the same size_t attrs_idx = 0; declaration is reused for hipMemcpyBatchAsync(..., &attr, &attrs_idx, num_attrs, ...). Note the ROCm path may pass num_attrs = 0 on ROCm 7.2.1–7.2.3 and older runtimes, in which case attrIdxs is ignored and the pattern is harmless there — but on ROCm 7.13+ (num_attrs = 1) the same OOB read applies.
Suggested fix
Allocate a per-descriptor index array. All-zero indices preserve current semantics exactly (every descriptor uses attrs[0]):
// hoisted outside the chunk loop; reused across chunks
std::vector<size_t> attr_idxs;
...
attr_idxs.assign(static_cast<size_t>(cnt), 0);
CUresult result = batch_fn(..., static_cast<size_t>(cnt), &attr,
attr_idxs.data(), 1, &fail_idx,
static_cast<CUstream>(stream));(A small stack array or std::array sized to the chunk cap would also work if heap allocation in this path is a concern; the vector is the minimal-diff option.)
Related
I kept the Python PR minimal (no C++ changes) to avoid compile-matrix impact, as noted there. Happy to submit a PR for this side too if maintainers prefer — the change itself is small; the cost is mainly re-validating the build.
Before submitting a new issue...
Your current environment
Environment info
🐛 Describe the bug
Out-of-bounds
attrIdxsin the C++ batch memcpy path (cache_kernels.cu)Summary
The C++ batch-copy path in
csrc/libtorch_stable/cache_kernels.cupasses a singlesize_tas theattrIdxsarray tocuMemcpyBatchAsync/hipMemcpyBatchAsync, regardless of the descriptor count — the same contract violation fixed for the Python path in #53860.The CUDA Driver API documentation for
cuMemcpyBatchAsyncspecifiesattrIdxsas an array ofcountelements (indices intoattrs, each< numAttrs). Passing&attrs_idxwhereattrs_idxis one stack scalar whilecountis large (typically thousands of 64 KiB KV blocks per offload) leaves the driver reading past the end of that scalar — an out-of-bounds read / undefined behavior that currently happens to work only because adjacent stack bytes happen to read as tolerable values.Location
Two sites, both in the batch-copy helper (around L130–L230 of
csrc/libtorch_stable/cache_kernels.cu):CUDA path (L169):
ROCm path (L196): the same
size_t attrs_idx = 0;declaration is reused forhipMemcpyBatchAsync(..., &attr, &attrs_idx, num_attrs, ...). Note the ROCm path may passnum_attrs = 0on ROCm 7.2.1–7.2.3 and older runtimes, in which caseattrIdxsis ignored and the pattern is harmless there — but on ROCm 7.13+ (num_attrs = 1) the same OOB read applies.Suggested fix
Allocate a per-descriptor index array. All-zero indices preserve current semantics exactly (every descriptor uses
attrs[0]):(A small stack array or
std::arraysized to the chunk cap would also work if heap allocation in this path is a concern; the vector is the minimal-diff option.)Related
vllm/v1/simple_kv_offload/cuda_mem_ops.py), including the API-contract analysis, a guard-page OOB probe, and a standalone verifier.cuMemcpyBatchAsyncsegfaults in the KV-offloading path (closed after "disappearing" on a nightly build — consistent with layout-dependent UB rather than a deterministic logic bug).I kept the Python PR minimal (no C++ changes) to avoid compile-matrix impact, as noted there. Happy to submit a PR for this side too if maintainers prefer — the change itself is small; the cost is mainly re-validating the build.
Before submitting a new issue...