[BUG] chunked_parquet_reader hangs or hits cudaErrorIllegalAddress on zstd DataPageV2 files: nvCOMP temp-size query is given the uncompressed level bytes

**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:

```cpp
temp_spans[i] = device_span(
page.page_data, static_cast(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**

```python
# 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_PRELOAD`ed 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.