CSIRTS // UNIFIED SECURITY ADVISORY FEEDSYS ● ONLINE · POWERED BY INTELFUSIONS.COM

GHSA-5jv2-g5wq-cmr4: vLLM: GGUF dequantize kernel int truncation exposes uninitialized GPU memory in multi-tenant serving

mediumCVSS 7.5CVE-2026-53923
Summary Integer truncation of tensor dimensions in vLLM's GGUF dequantize kernels (csrc/quantization/gguf/gguf_kernel.cu) causes partial tensor processing. The output tensor is allocated at full size via torch::empty (uninitialized memory), but the dequantize CUDA kernel processes only a truncated number of elements. The unfilled portion of the output tensor retains whatever was previously in GPU memory. In multi-tenant inference deployments, this residual GPU memory may contain tensor data from other users' inference requests, constituting information disclosure. Root Cause The to_cuda_ggml_t function pointer type at ggml-common.h:1067 declares its element count parameter as int (32-bit): using to_cuda_ggml_t = void (*)(const void * restrict x, dst_t * restrict y, int k, // 32-bit cudaStream_t stream); All dequantize kernel functions (dequantize_block_cuda, dequantize_row_q2_K_cuda, etc. in dequantize.cuh) inherit this int k parameter and use it as the kernel launch grid size: static void dequantize_block_cuda(..., const int k, cudaStream_t stream) { const int num_blocks = (k + 2*CUDA_DEQUANTIZE_BLOCK_SIZE - 1) / (2*CUDA_DEQUANTIZE_BLOCK_SIZE); dequantize_block<<<num_blocks, CUDA_DEQUANTIZE_BLOCK_SIZE, 0, stream>>>(vx, y, k); } In ggml_dequantize() at gguf_kernel.cu:85, the caller passes m * n (an int64_t product) to this int k parameter: at::Tensor DW = torch::empty({m, n}, options); // line 80: full-size, UNINITIALIZED // ... to_cuda((void*)W.data_ptr(), (scalar_t*)DW.data_ptr(), m * n, stream); // line 85: m*n truncated to int When m * n > INT_MAX, the truncated k is smaller than the actual tensor size. The kernel processes k elements. The remaining (m * n) - k elements in DW are never written and contain stale GPU memory. This is a single root cause -- the int type on the k parameter in to_cuda_ggml_t -- with a single fix: change int k to int64_t k. All dequantize functions inherit this type through the same typedef. Affected Functions All in csrc/q

Details

Source
GitHub Security Advisories (INTL · database · site)
Severity
medium — CVSS 7.5
Published
2026-06-17
Last updated
2026-07-17
Exploitation
Not in CISA KEV at last sync

Original advisory: https://github.com/advisories/GHSA-5jv2-g5wq-cmr4

Exploitation outlook

EPSS (FIRST.org) estimates each CVE’s probability of exploitation in the next 30 days — here is the CSIRTS.com read on those numbers.

Referenced CVEs

CVECSIRTS overviewExternal
CVE-2026-53923coverage & exploitation statusNVD · CVE.org

Same CVEs, other sources

How other CERTs, PSIRTs and databases cover the vulnerabilities in this advisory.

More from GitHub Security Advisories