[BUG] Invalid memory read in cudf::io::parquet::gpu::gpuDecodePageData

[davidwendt]

The nightly memcheck tests found an invalid global device memory read in cudf::io::parquet::gpu::gpuDecodePageData

========= Invalid __global__ read of size 4 bytes
=========     at 0x4dd0 in void cudf::io::parquet::gpu::<unnamed>::gpuDecodePageData<(int)256, unsigned char>(cudf::io::parquet::gpu::PageInfo *, cudf::device_span<const cudf::io::parquet::gpu::ColumnChunkDesc, (unsigned long)18446744073709551615>, unsigned long, unsigned long)
=========     by thread (81,0,0) in block (0,0,0)
=========     Address 0x7f0511e00ad8 is out of bounds
=========     and is inside the nearest allocation at 0x7f0511e00a00 of size 218 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x304e32]
=========                in /usr/lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame: [0x1488c]
=========                in /opt/conda/envs/test/bin/gtests/libcudf/../../../lib/libcudart.so.11.0
=========     Host Frame:cudaLaunchKernel [0x6c318]
=========                in /opt/conda/envs/test/bin/gtests/libcudf/../../../lib/libcudart.so.11.0
=========     Host Frame:cudf::io::parquet::gpu::DecodePageData(cudf::detail::hostdevice_vector<cudf::io::parquet::gpu::PageInfo>&, cudf::detail::hostdevice_vector<cudf::io::parquet::gpu::ColumnChunkDesc> const&, unsigned long, unsigned long, int, rmm::cuda_stream_view) [0x155aaef]
=========                in /opt/conda/envs/test/bin/gtests/libcudf/../../../lib/libcudf.so
=========     Host Frame:cudf::io::detail::parquet::reader::impl::decode_page_data(unsigned long, unsigned long) [0x1564ba9]
=========                in /opt/conda/envs/test/bin/gtests/libcudf/../../../lib/libcudf.so
=========     Host Frame:cudf::io::detail::parquet::reader::impl::read_chunk_internal(bool) [0x1566b2d]
=========                in /opt/conda/envs/test/bin/gtests/libcudf/../../../lib/libcudf.so
=========     Host Frame:cudf::io::detail::parquet::reader::impl::read(long, std::optional<int> const&, bool, cudf::host_span<std::vector<int, std::allocator<int> > const, 18446744073709551615ul>) [0x15672dc]
=========                in /opt/conda/envs/test/bin/gtests/libcudf/../../../lib/libcudf.so
=========     Host Frame:cudf::io::detail::parquet::reader::read(cudf::io::parquet_reader_options const&) [0x155d03b]
=========                in /opt/conda/envs/test/bin/gtests/libcudf/../../../lib/libcudf.so
=========     Host Frame:cudf::io::read_parquet(cudf::io::parquet_reader_options const&, rmm::mr::device_memory_resource*) [0x147505e]
=========                in /opt/conda/envs/test/bin/gtests/libcudf/../../../lib/libcudf.so

The error can be recreated using the PARQUET_TEST as follows:

compute-sanitizer --tool memcheck gtests/PARQUET_TEST --gtest_filter=ParquetWriterNumericTypeTest/2.SingleColumnWithNulls --rmm_mode=cuda