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

[davidwendt]

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():

if (_smem_stage != nullptr) {
  auto* const smem_stage = static_cast<uint8_t const*>(cuda::std::assume_aligned<16>(_smem_stage));
  auto const len = static_cast<int>(cuda::std::distance(_start, _end));
  if (len > 0 && len <= stage_capacity) {
    cuda::memcpy_async(group, _smem_stage, _start, static_cast<size_t>(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:

• Optimize RLE decoding by staging in shared memory #23090
• Optimize RLE decoding using a warp-balanced chunking approach #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.

[davidwendt]

@vyasr