Describe the bug
cudf::io::chunked_parquet_reader with a non-zero pass_read_limit hangs in a GPU kernel or aborts with cudaErrorIllegalAddress on zstd-compressed Parquet files whose data pages are DataPageV2 with non-empty repetition/definition level bytes (any nullable column written with data_page_version="2.0"). read_parquet on the same file returns the correct rows, and so does the chunked reader with pass_read_limit == 0 or with LIBCUDF_NVCOMP_POLICY=OFF.
When it does not hang, the first visible symptom is the warning
[warning] batched_decompress_get_temp_size_sync failed, falling back to batched_decompress_temp_size
followed by a sticky CUDA error on the next allocation:
MemoryError: std::bad_alloc: CUDA error (failed to allocate 16 bytes) at: .../rmm/cpp/src/mr/cuda_memory_resource.cpp:26: cudaErrorIllegalAddress an illegal memory access was encountered
compute-sanitizer places the illegal access in nvCOMP's zstd frame parser, launched from nvcompBatchedZstdDecompressGetTempSizeSync while the reader sizes its decompression scratch for the subpass:
========= Invalid __global__ read of size 1 bytes
========= at void zstd::gather_frame_blocks<(unsigned long)64>(unsigned long, const unsigned char *const *, unsigned long *, unsigned long *, nvcompStatus_t *)+0x390
========= by thread (1,0,0) in block (0,0,0)
========= Access to 0xfea5d368318f is out of bounds
========= and is 406,920 bytes after the nearest allocation at 0xfea5d361fc00 of size 8 bytes
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame: nvcompBatchedZstdDecompressGetTempSizeSync in libnvcomp.so.5
========= Host Frame: cudf::io::detail::nvcomp::(anonymous namespace)::batched_decompress_temp_size_ex(...) in libcudf.so
========= Host Frame: cudf::io::detail::nvcomp::batched_decompress_temp_size_ex(...) in libcudf.so
========= Host Frame: cudf::io::detail::get_decompression_scratch_size_ex(...) in libcudf.so
========= Host Frame: cudf::io::parquet::detail::compute_decompression_scratch_sizes(...) in libcudf.so
========= Host Frame: cudf::io::parquet::detail::reader_impl::setup_next_subpass(cudf::io::parquet::detail::reader_impl::read_mode) in libcudf.so
========= Host Frame: cudf::io::parquet::detail::reader_impl::has_next() in libcudf.so
========= Host Frame: cudf::io::chunked_parquet_reader::has_next() const in libcudf.so
Likely cause
compute_decompression_scratch_sizes (cpp/src/io/parquet/reader_impl_chunking_utils.cu, lines 767-768 on release/26.08, 772-773 on main @ 60436a8) hands nvCOMP the raw page span of every page of the codec:
temp_spans[i] = device_span<uint8_t const>(
page.page_data, static_cast<size_t>(page.compressed_page_size));
For a DataPageV2 page page_data starts with the uncompressed repetition/definition level bytes, and the page may also carry is_compressed == false. The decompression path in the same file (set_parameters, lines ~537-550) handles both: it skips page.lvl_bytes[DEFINITION] + page.lvl_bytes[REPETITION] and only submits pages with is_compressed set. The scratch-size path does neither, so nvcompBatchedZstdDecompressGetTempSizeSync (nvCOMP >= 5.0, introduced by #19616) parses the level bytes as a zstd frame header on the device. Depending on what those bytes happen to say, the query returns a failure status (the warning above, then a silent fallback), reads out of bounds (cudaErrorIllegalAddress), or walks "blocks" forever (a kernel that never terminates: the process sits at 96-100 % GPU utilisation until killed).
That also explains what does and does not trigger it: DataPageV1 files, non-nullable columns (zero level bytes) and snappy never print the warning; pass_read_limit == 0 skips compute_decompression_scratch_sizes entirely, which is why read_parquet is unaffected.
Steps/Code to reproduce bug
# pyarrow 21.0.0 writer + pylibcudf 26.08.00 reader
import os, numpy as np, pyarrow as pa, pyarrow.parquet as pq
from pylibcudf.io.parquet import ChunkedParquetReader, ParquetReaderOptions, read_parquet
from pylibcudf.io.types import SourceInfo
n = 200_000
rng = np.random.default_rng(int(os.environ.get("SEED", "2")))
nation = rng.integers(0, 25, n, dtype=np.int64) # 25 distinct values -> RLE_DICTIONARY
null_frac = float(os.environ.get("NULL_FRAC", "0.1")) # outcome depends on the level bytes, see below
tbl = pa.table({
"key": np.arange(n, dtype=np.int64),
"nation": pa.array(nation, mask=rng.random(n) < null_frac),
})
pq.write_table(tbl, "repro.parquet", compression="zstd", compression_level=3,
data_page_version="2.0", write_page_index=True, write_statistics=False)
opts = ParquetReaderOptions.builder(SourceInfo(["repro.parquet"])).build()
opts.set_column_names(["nation"])
print("read_parquet rows:", read_parquet(opts).tbl.num_rows(), flush=True) # works
reader = ChunkedParquetReader(opts, chunk_read_limit=0, pass_read_limit=1024_000_000)
rows = 0
while reader.has_next(): # hangs / crashes here
rows += reader.read_chunk().tbl.num_rows()
print("chunked rows:", rows)Output with the defaults (SEED=2 NULL_FRAC=0.1): the process prints read_parquet rows: 200000 and never returns from has_next(); nvidia-smi shows 96 % GPU utilisation until the process is killed. Same with the unmodified libcudf-26.08.00-cuda13_260805_ff5b362d conda package (LD_PRELOADed to be sure). LIBCUDF_NVCOMP_POLICY=OFF on the same file prints chunked rows: 200000.
Because nvCOMP is fed the level bytes, the outcome depends on their content (same code, same shape, different random draw):
SEED / NULL_FRAC |
result |
| 2 / 0.1, 3 / 0.5 |
hang (both libcudf builds I have) |
| 5 / 0.5, 6 / 0.5 |
cudaErrorIllegalAddress after the warning, or only the warning, depending on the build/allocation layout |
| 3, 4, 5 / 0.1 |
warning only, rows correct |
| 1 / 0.1, 1 / 0.5, 6 / 0.1 |
no warning, rows correct |
Other shapes I tried (all pyarrow, data_page_version="2.0", write_page_index=True, write_statistics=False): a 9-column TPC-H customer-like table, 1.5M rows, zstd 19, 8 MiB data pages, 128 MiB row groups (3 row groups): selecting the dictionary-encoded INT64 column crashes every time with cudaErrorIllegalAddress, selecting a 44 MiB random-text string column hangs, selecting all columns hangs; zstd level 1/3/19, 1 MiB/8 MiB pages, 50k-1.5M rows behave the same way. pass_read_limit 1 and 1e6 happened to survive with two warnings, 1.024e9 crashed. read_parquet on every file: correct rows. data_page_version="1.0", non-nullable schema fields, compression="snappy": no warning, correct rows.
The same failure was first hit on a GB200 node (aarch64, CUDA 13.0, libcudf 26.08.00 conda package, nvCOMP 5.3.0.16) from C++ (cudf::io::chunked_parquet_reader constructed with pre-parsed footers) on TPC-H customer files written by parquet-rs (zstd level 19, DataPageV2, offset index only): identical warning and cudaErrorIllegalAddress from has_next(); LIBCUDF_NVCOMP_POLICY=OFF fixed it there too.
Expected behavior
The chunked reader returns the same rows as read_parquet. The scratch-size estimate should be computed on the spans that are actually decompressed (skip the V2 level bytes and pages with is_compressed == false); a failed GetTempSizeSync should not be able to leave a sticky device error or an unbounded kernel behind.
Environment overview
- Environment location: bare metal, NVIDIA DGX Spark (GB10, sm_121, driver 580.95.05), Ubuntu 24.04.3, aarch64; first seen on NVIDIA GB200 (sm_100, CUDA 13.0)
- Method of cuDF install: conda (
libcudf 26.08.00 cuda13_260805_ff5b362d, libnvcomp 5.3.0.16, librmm 26.08.00, cuda-version 13.3, pylibcudf/cudf 26.08.00 cuda13_cp311_abi3_260805_ff5b362d, Python 3.11.16, pyarrow 21.0.0, numpy 2.4.6)
main @ 60436a8 carries the same code in compute_decompression_scratch_sizes, so I expect it to reproduce there as well (not run).
Environment details
print_env.sh (abridged)
***OS Information***
PRETTY_NAME="Ubuntu 24.04.3 LTS"
VERSION="24.04.3 LTS (Noble Numbat)"
Linux dgx-spark 6.14.0-1013-nvidia #13-Ubuntu SMP PREEMPT_DYNAMIC Wed Oct 29 06:01:19 UTC 2025 aarch64 aarch64 aarch64 GNU/Linux
***GPU Information***
| NVIDIA-SMI 580.95.05 Driver Version: 580.95.05 CUDA Version: 13.0 |
| 0 NVIDIA GB10 On | 0000000F:01:00.0 Off | N/A |
***conda packages***
cuda-cudart 13.3.29 h8f3c8d4_0
cuda-sanitizer-api 13.3.75 h299c5c6_0
cuda-version 13.3 hcbadf70_3
cudf 26.08.00 cuda13_cp311_abi3_260805_ff5b362d
libcudf 26.08.00 cuda13_260805_ff5b362d
libkvikio 26.08.00 cuda13_260805_5a77056e
libnvcomp 5.3.0.16 he387df4_0
libnvjitlink 13.3.33 h8f3c8d4_0
librmm 26.08.00 cuda13_260805_42d059f1
numpy 2.4.6 py311hecca567_0
pyarrow 21.0.0 py311hfecb2dc_3
pylibcudf 26.08.00 cuda13_cp311_abi3_260805_ff5b362d
python 3.11.16 ha505bbe_0_cpython
rapids-logger 0.2.3 h4f43097_0
rmm 26.08.00 cuda13_cp311_abi3_260805_42d059f1
Additional context
compute-sanitizer from CUDA 13.0 reports the first failure as an "Invalid shared read" inside a CUB DeviceScanKernel launched from decode_page_headers (on both the GB200 and the GB10); the CUDA 13.3 sanitizer that ships in the conda env (cuda-sanitizer-api) points at the nvCOMP kernel above. The 13.0 report looks like a tool/CCCL mismatch; mentioned in case someone else chases the CUB scan first.
Describe the bug
cudf::io::chunked_parquet_readerwith a non-zeropass_read_limithangs in a GPU kernel or aborts withcudaErrorIllegalAddresson zstd-compressed Parquet files whose data pages are DataPageV2 with non-empty repetition/definition level bytes (any nullable column written withdata_page_version="2.0").read_parqueton the same file returns the correct rows, and so does the chunked reader withpass_read_limit == 0or withLIBCUDF_NVCOMP_POLICY=OFF.When it does not hang, the first visible symptom is the warning
followed by a sticky CUDA error on the next allocation:
compute-sanitizerplaces the illegal access in nvCOMP's zstd frame parser, launched fromnvcompBatchedZstdDecompressGetTempSizeSyncwhile the reader sizes its decompression scratch for the subpass:Likely cause
compute_decompression_scratch_sizes(cpp/src/io/parquet/reader_impl_chunking_utils.cu, lines 767-768 onrelease/26.08, 772-773 onmain@ 60436a8) hands nvCOMP the raw page span of every page of the codec:For a DataPageV2 page
page_datastarts with the uncompressed repetition/definition level bytes, and the page may also carryis_compressed == false. The decompression path in the same file (set_parameters, lines ~537-550) handles both: it skipspage.lvl_bytes[DEFINITION] + page.lvl_bytes[REPETITION]and only submits pages withis_compressedset. The scratch-size path does neither, sonvcompBatchedZstdDecompressGetTempSizeSync(nvCOMP >= 5.0, introduced by #19616) parses the level bytes as a zstd frame header on the device. Depending on what those bytes happen to say, the query returns a failure status (the warning above, then a silent fallback), reads out of bounds (cudaErrorIllegalAddress), or walks "blocks" forever (a kernel that never terminates: the process sits at 96-100 % GPU utilisation until killed).That also explains what does and does not trigger it: DataPageV1 files, non-nullable columns (zero level bytes) and snappy never print the warning;
pass_read_limit == 0skipscompute_decompression_scratch_sizesentirely, which is whyread_parquetis unaffected.Steps/Code to reproduce bug
Output with the defaults (
SEED=2 NULL_FRAC=0.1): the process printsread_parquet rows: 200000and never returns fromhas_next();nvidia-smishows 96 % GPU utilisation until the process is killed. Same with the unmodifiedlibcudf-26.08.00-cuda13_260805_ff5b362dconda package (LD_PRELOADed to be sure).LIBCUDF_NVCOMP_POLICY=OFFon the same file printschunked rows: 200000.Because nvCOMP is fed the level bytes, the outcome depends on their content (same code, same shape, different random draw):
SEED/NULL_FRACcudaErrorIllegalAddressafter the warning, or only the warning, depending on the build/allocation layoutOther shapes I tried (all pyarrow,
data_page_version="2.0",write_page_index=True,write_statistics=False): a 9-column TPC-Hcustomer-like table, 1.5M rows, zstd 19, 8 MiB data pages, 128 MiB row groups (3 row groups): selecting the dictionary-encoded INT64 column crashes every time withcudaErrorIllegalAddress, selecting a 44 MiB random-text string column hangs, selecting all columns hangs; zstd level 1/3/19, 1 MiB/8 MiB pages, 50k-1.5M rows behave the same way.pass_read_limit1 and 1e6 happened to survive with two warnings, 1.024e9 crashed.read_parqueton every file: correct rows.data_page_version="1.0", non-nullable schema fields,compression="snappy": no warning, correct rows.The same failure was first hit on a GB200 node (aarch64, CUDA 13.0, libcudf 26.08.00 conda package, nvCOMP 5.3.0.16) from C++ (
cudf::io::chunked_parquet_readerconstructed with pre-parsed footers) on TPC-Hcustomerfiles written by parquet-rs (zstd level 19, DataPageV2, offset index only): identical warning andcudaErrorIllegalAddressfromhas_next();LIBCUDF_NVCOMP_POLICY=OFFfixed it there too.Expected behavior
The chunked reader returns the same rows as
read_parquet. The scratch-size estimate should be computed on the spans that are actually decompressed (skip the V2 level bytes and pages withis_compressed == false); a failedGetTempSizeSyncshould not be able to leave a sticky device error or an unbounded kernel behind.Environment overview
libcudf 26.08.00 cuda13_260805_ff5b362d,libnvcomp 5.3.0.16,librmm 26.08.00,cuda-version 13.3,pylibcudf/cudf26.08.00cuda13_cp311_abi3_260805_ff5b362d, Python 3.11.16, pyarrow 21.0.0, numpy 2.4.6)main@ 60436a8 carries the same code incompute_decompression_scratch_sizes, so I expect it to reproduce there as well (not run).Environment details
print_env.sh (abridged)
Additional context
compute-sanitizerfrom CUDA 13.0 reports the first failure as an "Invalid shared read" inside a CUBDeviceScanKernellaunched fromdecode_page_headers(on both the GB200 and the GB10); the CUDA 13.3 sanitizer that ships in the conda env (cuda-sanitizer-api) points at the nvCOMP kernel above. The 13.0 report looks like a tool/CCCL mismatch; mentioned in case someone else chases the CUB scan first.