[BUG] Debug builds crash with memcpy: source range is invalid in preprocess_levels_kernel (RLE level-stream shared-memory staging)

**Describe the bug**

In **Debug** builds only, several parquet reader tests crash with a device-side assert:
```
libcudacxx/include/cuda/std/__cstring/memcpy.h:49: memcpy: block: [N,0,0], thread: [M,0,0]
Assertion memcpy: source range is invalid failed.
CUDA Error detected. cudaErrorAssert device-side assert triggered
```
The crash happens in `preprocess_levels_kernel` (`cpp/src/io/parquet/decode_preprocess.cu`), which decodes definition/repetition levels via `rle_stream::decode_next_chunked()` in `cpp/src/io/parquet/rle_stream.cuh`.

Bisected this to the shared-memory staging path added in `rle_stream::init()`:

```cpp
if (_smem_stage != nullptr) {
auto* const smem_stage = static_cast(cuda::std::assume_aligned<16>(_smem_stage));
auto const len = static_cast(cuda::std::distance(_start, _end));
if (len > 0 && len <= stage_capacity) {
cuda::memcpy_async(group, _smem_stage, _start, static_cast(len), *_copy_barrier);
cur = smem_stage;
end = smem_stage + len;
}
}
```

Passing `nullptr` instead of `stage` for both the REPETITION and DEFINITION decoder `init()` calls in
`decode_preprocess.cu` (i.e. disabling shared-memory staging entirely and always decoding from global memory) makes all failing tests pass. Re-enabling staging for either decoder reproduces the crash. I was not able to pin the exact defective line further — `cuda-gdb` crashes with its own internal error.
Using `compute-sanitizer --show-backtrace=yes` shows host-side stacks and SASS/line-table correlation which may help.

Bisected to the combination of:
- #23090
- #23271

Both touch `cpp/src/io/parquet/rle_stream.cuh` / `cpp/src/io/parquet/decode_preprocess.cu`.

**Steps/Code to reproduce bug**

Build `libcudf` in **Debug** mode (`CMAKE_BUILD_TYPE=Debug`), then run any of:

```
UTILITIES_TEST --gtest_filter=PinnedMemoryTest.MemoryResourceGetAndSet
PARQUET_TEST --gtest_filter=ParquetChunkedReaderTest.TestChunkedReadSimpleData
HYBRID_SCAN_TEST --gtest_filter=HybridScanMultifileTest.MaterializeLists
LARGE_STRINGS_TEST --gtest_filter=ParquetStringsTest.ChunkedReadNestedLargeStrings
STREAM_IO_PARQUET_TEST --gtest_filter=ParquetTest.ChunkedOperations
```

All five fail identically with the assert above, and all five pass when shared-memory staging is disabled per the bisection described.

Start with cpp/src/io/parquet/rle_stream.cuh and cpp/src/io/parquet/decode_preprocess.cu, focusing on the shared-memory staging path in rle_stream::init() and the REPETITION and DEFINITION decoder init() calls. Build libcudf in Debug mode and run a listed failing filter, using compute-sanitizer --show-backtrace=yes as suggested. Done means the staging path no longer triggers the invalid memcpy assertion and the affected parquet tests pass with staging enabled.

Written by the indexing model from the issue text.