Repository navigation
Uninitialized device-memory read in map_open_close #6
Description
Uninitialized device-memory read in map_open_close
Environment
- cuJSON commit:
2ac7d3dcd7ad1ff64ebdb14022bf94c59b3b4953(branchmaster) - OS: Ubuntu 22.04.5 LTS
- GPU: NVIDIA A40 (Ampere,
sm_86) - NVIDIA driver: 590.48.01
- CUDA toolkit: 13.1 (
nvccV13.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.binSanitizer 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.cuharness_afl_main.cpptrigger_cuJSONLinesIterator.bininitcheck.log
Activity
Merge pull request #17 from AutomataLab/fix/parser-uninit-and-mergech…
ashkanvg commentedon Sep 10, 2026
Closing as a false positive — this is not a cuJSON bug. Fixed/explained in #17.
Root cause. The reported uninitialized reads all originate inside
thrust::stable_sort_by_keyover azip_iteratorwhose value type istuple<uint32_t, uint8_t>. That type is 8 bytes wide for 5 bytes of payload, so it carries 3 bytes of struct padding. Thrust materialises a temporary tuple array by assigning only the two members, then CUB'sBlockLoadreads whole 8-byte tuples — including padding it never wrote.Proof. A standalone program containing no cuJSON code at all, in which every input byte is explicitly
cudaMemset:The padding is never compared by the radix sort and never scattered back to the output arrays, so it cannot influence results.
Real defects found while investigating this, fixed in #17:
mergeChunksallocatedsizeof(uint32_t)*N*ROW2 + 3=8N+3bytes but writes through element2N+1, needing8N+8— a genuine 5-byte heap overflow on every non-empty JSON Lines merge, confirmed with a guard canary. ASan could not see it because the buffer iscudaMallocHost.pair_posrow was copied to the host partly uninitialized and published throughcuJSONResult. Now zeroed.After #17, cuJSON-attributable initcheck findings are zero on both paths (standard 22 → 20, JSON Lines 60 → 53); every remaining report routes through
sort.inl:125.Fix mergeChunks buffer overflow and uninitialized pair_pos copy