Repository navigation
Uninitialized device-memory read passed into CUB radix sort #7
Description
Uninitialized device-memory read passed into CUB radix sort
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 — sort integration (parse_standard_json.cu:1584-1611, parse_json_lines.cu:1070-1093). Reproduces across cuJSONIterator, cuJSONLinesIterator, parse_json_lines, and parse_standard_json.
Severity
Medium (CWE-908, Use of Uninitialized Resource).
Description
cuJSON passes partially initialized key/value buffers into the thrust/CUB sorting stage. cub::DeviceRadixSortSingleTileKernel then loads uninitialized tail bytes while sorting structural positions.
Root cause
The buffers handed to the radix sort are sized for a padded element count, but only the populated structural/pair elements are initialized. CUB sorts the full allocated range, so it loads the uninitialized tail. This reproduces across all four entry points because they share the same sort integration.
Call path:
harness LLVMFuzzerTestOneInput
-> parse_standard_json / parse_json_lines / iterator entry points
-> CUB DeviceRadixSortSingleTileKernel (uninitialized read of sort buffer tail)
Impact
Systemic parse-tree corruption across standard JSON, JSON Lines, and iterator entry points. Malformed inputs can produce nondeterministic parse results and sanitizer-visible device reads.
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_cuJSONIterator.cu harness_afl_main.cpp \
--compiler-bindir g++-13 -L${CUDA_HOME}/lib64 -lcudart \
-o /tmp/cujson_cub_sort_initcheck.bin
compute-sanitizer --tool=initcheck --report-api-errors=no --error-exitcode 99 \
/tmp/cujson_cub_sort_initcheck.bin trigger_cuJSONIterator.binSanitizer evidence
========= Uninitialized __global__ memory read of size 4 bytes
========= Uninitialized __global__ memory read of size 1 bytes (x6)
========= Uninitialized access to 0x7f3cf6400480 on access by cudaMemcpy source
========= Uninitialized access to 0x1938000420 on access by cudaMemcpy source
========= ERROR SUMMARY: 9 errors
Suggested fix
Track the exact initialized element count passed to the sorting kernels, zero-fill padded storage, and avoid sorting beyond the populated structural/pair arrays.
Attachments
The following files from this finding are attached:
llm_harness_cuJSONIterator.cuharness_afl_main.cpptrigger_cuJSONIterator.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