Uninitialized device-memory read in `map_open_close`

[TarekIbnZiad]

Uninitialized device-memory read in map_open_close

Environment

• cuJSON commit: 2ac7d3dcd7ad1ff64ebdb14022bf94c59b3b4953 (branch master)
• OS: Ubuntu 22.04.5 LTS
• GPU: NVIDIA A40 (Ampere, sm_86)
• NVIDIA driver: 590.48.01
• CUDA toolkit: 13.1 (nvcc V13.1.115)

Component

cuJSON — JSON Lines path, map_open_close kernel (parse_json_lines.cu:979-992, parse_json_lines.cu:1064-1093). Affects cuJSONLinesIterator and parse_json_lines.

Severity

Medium (CWE-908, Use of Uninitialized Resource).

Description

map_open_close reads 32-bit words from open_close_bitmap. For JSON Lines inputs whose structural count does not fill a full word, the final padded word contains bytes that were allocated but never initialized, so the kernel reads uninitialized device memory.

Root cause

The open_close_bitmap allocation is padded to a word boundary, but the padding bytes in the final word are not initialized when the structural count is not word-aligned. map_open_close then reads the full word, including those uninitialized bytes.

Call path:

harness LLVMFuzzerTestOneInput
  -> cuJSONLinesIterator, parse_json_lines
  -> map_open_close   (uninitialized read)

Impact

Uninitialized structural-map bits can corrupt bracket pairing or depth computation, producing incorrect parse results or later GPU memory errors.

Reproduction

CUDA_HOME=${CUDA_HOME:-/usr/local/cuda}
nvcc -std=c++17 -arch=native -O2 -g -lineinfo \
  -I<cujson-src> -I${CUDA_HOME}/include \
  llm_harness_cuJSONLinesIterator.cu harness_afl_main.cpp \
  --compiler-bindir g++-13 -L${CUDA_HOME}/lib64 -lcudart \
  -o /tmp/cujson_map_open_close_initcheck.bin

compute-sanitizer --tool=initcheck --report-api-errors=no --error-exitcode 99 \
  /tmp/cujson_map_open_close_initcheck.bin trigger_cuJSONLinesIterator.bin

Sanitizer evidence

========= Uninitialized __global__ memory read of size 4 bytes
========= Uninitialized __global__ memory read of size 1 bytes  (x6)
=========   Uninitialized access to 0x7fa1e2400480 on access by cudaMemcpy source
=========   Uninitialized access to 0x1938000480 on access by cudaMemcpy source
=========   Uninitialized access to 0x19380004a0 on access by cudaMemcpy source
=========   Uninitialized access to 0x19380004c0 on access by cudaMemcpy source
=========   Uninitialized access to 0x19380004e0 on access by cudaMemcpy source

Suggested fix

Zero-fill the padded open_close_bitmap allocation, and clamp map_open_close reads to the initialized bytes in the final word.

Attachments

The following files from this finding are attached:

• llm_harness_cuJSONLinesIterator.cu
• harness_afl_main.cpp
• trigger_cuJSONLinesIterator.bin
• initcheck.log

Archive.zip