Problem
cudf::concatenate treats every DICTIONARY32 indices child as INT32. With
INT8 indices this causes an illegal memory access; with INT16 it overruns
the output buffer and compute-sanitizer reports out-of-bounds accesses. The
error is asynchronous and may surface in a later CUDA call. INT32 indices
are unaffected.
Reproducer
auto dict = cudf::dictionary::encode(
plain_column, cudf::data_type{cudf::type_id::INT8}); // INT16 is also affected
auto out = cudf::concatenate(std::vector<cudf::column_view>{*dict, *dict});
cudaDeviceSynchronize();On cuDF 26.08 and current main at d45003e0d0, the INT8 case reports
cudaErrorIllegalAddress; the INT16 case reports out-of-bounds reads and
writes under compute-sanitizer.
Cause
map_indices_fn reads with element<size_type>(), and the remap writes with
begin<size_type>(), even when the allocated indices column uses one- or
two-byte elements. This was introduced by #22807.
Expected behavior
- Read and write indices using their actual integral type.
- For mixed input widths, use the widest input type.
- Widen further when the merged keys no longer fit that type, for example from
INT8 to INT16 when the result has 200 keys.
- Preserve decoded values and nulls.
Related work
Problem
cudf::concatenatetreats everyDICTIONARY32indices child asINT32. WithINT8indices this causes an illegal memory access; withINT16it overrunsthe output buffer and compute-sanitizer reports out-of-bounds accesses. The
error is asynchronous and may surface in a later CUDA call.
INT32indicesare unaffected.
Reproducer
On cuDF 26.08 and current
mainatd45003e0d0, theINT8case reportscudaErrorIllegalAddress; theINT16case reports out-of-bounds reads andwrites under compute-sanitizer.
Cause
map_indices_fnreads withelement<size_type>(), and the remap writes withbegin<size_type>(), even when the allocated indices column uses one- ortwo-byte elements. This was introduced by #22807.
Expected behavior
INT8toINT16when the result has 200 keys.Related work