Skip to content

refactor(parquet): introduce string_size_scan_state for string size kernels #23496

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
rapids-bot merged 4 commits into from
Aug 10, 2026
Merged

refactor(parquet): introduce string_size_scan_state for string size kernels #23496

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
ad71ecf
refactor(parquet): extract page_decode_output_state and narrow string…
vyasr Jul 29, 2026
2baed40
Update struct description
vyasr Jul 31, 2026
250bb31
Merge remote-tracking branch 'upstream/main' into parquet-substruct-r…
vyasr Aug 10, 2026
17b1a06
refactor(parquet): drop leftover string_offset_scan_state size static…
vyasr Aug 10, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
37 changes: 19 additions & 18 deletions cpp/src/io/parquet/decode_fixed.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -159,7 +159,7 @@ __device__ void decode_fixed_width_values(
auto const data_out = s->nesting_info[leaf_level_index].data_out;

Type const dtype = s->setup.col.physical_type;
uint32_t const dtype_len = s->dtype_len;
uint32_t const dtype_len = s->output_cvt.dtype_len;

int const skipped_leaf_values = s->setup.page.skipped_leaf_values;

Expand Down Expand Up @@ -200,9 +200,9 @@ __device__ void decode_fixed_width_values(
read_fixed_width_value_fast(s, sb, src_pos, static_cast<uint2*>(dst));
break;
default:
if (s->dtype_len_in <= sizeof(int32_t)) {
if (s->output_cvt.dtype_len_in <= sizeof(int32_t)) {
read_fixed_width_byte_array_as_int(s, sb, src_pos, static_cast<int32_t*>(dst));
} else if (s->dtype_len_in <= sizeof(int64_t)) {
} else if (s->output_cvt.dtype_len_in <= sizeof(int64_t)) {
read_fixed_width_byte_array_as_int(s, sb, src_pos, static_cast<int64_t*>(dst));
} else {
read_fixed_width_byte_array_as_int(s, sb, src_pos, static_cast<__int128_t*>(dst));
Expand All @@ -214,15 +214,15 @@ __device__ void decode_fixed_width_values(
} else if (dtype == Type::INT96) {
read_int96_timestamp(s, sb, src_pos, static_cast<int64_t*>(dst));
} else if (dtype_len == 8) {
if (s->dtype_len_in == 4) {
if (s->output_cvt.dtype_len_in == 4) {
// Reading INT32 TIME_MILLIS into 64-bit DURATION_MILLISECONDS
// TIME_MILLIS is the only duration type stored as int32:
// https://github.com/apache/parquet-format/blob/master/LogicalTypes.md#deprecated-time-convertedtype
auto const dst_ptr = static_cast<uint32_t*>(dst);
read_fixed_width_value_fast(s, sb, src_pos, dst_ptr);
// zero out most significant bytes
cuda::std::memset(dst_ptr + 1, 0, sizeof(int32_t));
} else if (s->ts_scale) {
} else if (s->output_cvt.ts_scale) {
read_int64_timestamp(s, sb, src_pos, static_cast<int64_t*>(dst));
} else {
read_fixed_width_value_fast(s, sb, src_pos, static_cast<uint2*>(dst));
Expand Down Expand Up @@ -253,12 +253,12 @@ __device__ inline void decode_fixed_width_split_values(
auto const data_len = cuda::std::distance(s->stream.data_start, s->stream.data_end);

// Check malformed BYTE_STREAM_SPLIT pages
if (s->dtype_len_in <= 0 or data_len <= 0) {
if (s->output_cvt.dtype_len_in <= 0 or data_len <= 0) {
if (t == 0) { s->set_error_code(decode_error::INVALID_BYTE_STREAM_SPLIT_SIZE); }
return;
}

auto const num_values = data_len / s->dtype_len_in;
auto const num_values = data_len / s->output_cvt.dtype_len_in;

int const skipped_leaf_values = s->setup.page.skipped_leaf_values;

Expand Down Expand Up @@ -288,7 +288,7 @@ __device__ inline void decode_fixed_width_split_values(
}
}();

uint32_t const dtype_len = s->dtype_len;
uint32_t const dtype_len = s->output_cvt.dtype_len;
uint8_t const* const src = s->stream.data_start + src_pos;
uint8_t* const dst = data_out + static_cast<size_t>(dst_pos) * dtype_len;
auto const is_decimal = s->setup.col.logical_type.has_value() and
Expand All @@ -300,17 +300,17 @@ __device__ inline void decode_fixed_width_split_values(
case Type::INT32: gpuOutputByteStreamSplit<int32_t>(dst, src, num_values); break;
case Type::INT64: gpuOutputByteStreamSplit<int64_t>(dst, src, num_values); break;
case Type::FIXED_LEN_BYTE_ARRAY:
if (s->dtype_len_in <= sizeof(int32_t)) {
if (s->output_cvt.dtype_len_in <= sizeof(int32_t)) {
gpuOutputSplitFixedLenByteArrayAsInt(
reinterpret_cast<int32_t*>(dst), src, num_values, s->dtype_len_in);
reinterpret_cast<int32_t*>(dst), src, num_values, s->output_cvt.dtype_len_in);
break;
} else if (s->dtype_len_in <= sizeof(int64_t)) {
} else if (s->output_cvt.dtype_len_in <= sizeof(int64_t)) {
gpuOutputSplitFixedLenByteArrayAsInt(
reinterpret_cast<int64_t*>(dst), src, num_values, s->dtype_len_in);
reinterpret_cast<int64_t*>(dst), src, num_values, s->output_cvt.dtype_len_in);
break;
} else if (s->dtype_len_in <= sizeof(__int128_t)) {
} else if (s->output_cvt.dtype_len_in <= sizeof(__int128_t)) {
gpuOutputSplitFixedLenByteArrayAsInt(
reinterpret_cast<__int128_t*>(dst), src, num_values, s->dtype_len_in);
reinterpret_cast<__int128_t*>(dst), src, num_values, s->output_cvt.dtype_len_in);
break;
}
// unsupported decimal precision
Expand All @@ -319,15 +319,16 @@ __device__ inline void decode_fixed_width_split_values(
default: s->set_error_code(decode_error::UNSUPPORTED_ENCODING);
}
} else if (dtype_len == 8) {
if (s->dtype_len_in == 4) {
if (s->output_cvt.dtype_len_in == 4) {
// Reading INT32 TIME_MILLIS into 64-bit DURATION_MILLISECONDS
// TIME_MILLIS is the only duration type stored as int32:
// https://github.com/apache/parquet-format/blob/master/LogicalTypes.md#deprecated-time-convertedtype
gpuOutputByteStreamSplit<int32_t>(dst, src, num_values);
// zero out most significant bytes
cuda::std::memset(dst + sizeof(int32_t), 0, sizeof(int32_t));
} else if (s->ts_scale) {
gpuOutputSplitInt64Timestamp(reinterpret_cast<int64_t*>(dst), src, num_values, s->ts_scale);
} else if (s->output_cvt.ts_scale) {
gpuOutputSplitInt64Timestamp(
reinterpret_cast<int64_t*>(dst), src, num_values, s->output_cvt.ts_scale);
} else {
gpuOutputByteStreamSplit<int64_t>(dst, src, num_values);
}
Expand Down Expand Up @@ -1306,7 +1307,7 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size_t, 8)
uint32_t const dtype_len = [&]() -> uint32_t {
if constexpr (is_dict_int32_t) { return sizeof(int32_t); }
if constexpr (has_strings_t) { return sizeof(cudf::size_type); }
return s->dtype_len;
return s->output_cvt.dtype_len;
}();
int const num_values = [&]() {
if constexpr (has_lists_t) {
Expand Down
48 changes: 27 additions & 21 deletions cpp/src/io/parquet/page_data.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -82,10 +82,10 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
bool const process_nulls = should_process_nulls(s);

auto const data_len = cuda::std::distance(s->stream.data_start, s->stream.data_end);
auto const num_values = data_len / s->dtype_len_in;
auto const num_values = data_len / s->output_cvt.dtype_len_in;

// Check malformed BYTE_STREAM_SPLIT pages
if (s->dtype_len_in <= 0 or data_len <= 0) {
if (s->output_cvt.dtype_len_in <= 0 or data_len <= 0) {
cg::invoke_one(block, [&]() {
set_error(static_cast<kernel_error::value_type>(decode_error::INVALID_BYTE_STREAM_SPLIT_SIZE),
error_code);
Expand Down Expand Up @@ -162,7 +162,7 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
// nesting level that is storing actual leaf values
int leaf_level_index = s->setup.col.max_nesting_depth - 1;

uint32_t dtype_len = s->dtype_len;
uint32_t dtype_len = s->output_cvt.dtype_len;
uint8_t const* src = s->stream.data_start + val_src_pos;
uint8_t* dst =
nesting_info_base[leaf_level_index].data_out + static_cast<size_t>(dst_pos) * dtype_len;
Expand All @@ -175,17 +175,17 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
case Type::INT32: gpuOutputByteStreamSplit<int32_t>(dst, src, num_values); break;
case Type::INT64: gpuOutputByteStreamSplit<int64_t>(dst, src, num_values); break;
case Type::FIXED_LEN_BYTE_ARRAY:
if (s->dtype_len_in <= sizeof(int32_t)) {
if (s->output_cvt.dtype_len_in <= sizeof(int32_t)) {
gpuOutputSplitFixedLenByteArrayAsInt(
reinterpret_cast<int32_t*>(dst), src, num_values, s->dtype_len_in);
reinterpret_cast<int32_t*>(dst), src, num_values, s->output_cvt.dtype_len_in);
break;
} else if (s->dtype_len_in <= sizeof(int64_t)) {
} else if (s->output_cvt.dtype_len_in <= sizeof(int64_t)) {
gpuOutputSplitFixedLenByteArrayAsInt(
reinterpret_cast<int64_t*>(dst), src, num_values, s->dtype_len_in);
reinterpret_cast<int64_t*>(dst), src, num_values, s->output_cvt.dtype_len_in);
break;
} else if (s->dtype_len_in <= sizeof(__int128_t)) {
} else if (s->output_cvt.dtype_len_in <= sizeof(__int128_t)) {
gpuOutputSplitFixedLenByteArrayAsInt(
reinterpret_cast<__int128_t*>(dst), src, num_values, s->dtype_len_in);
reinterpret_cast<__int128_t*>(dst), src, num_values, s->output_cvt.dtype_len_in);
break;
}
// unsupported decimal precision
Expand All @@ -194,16 +194,16 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
default: s->set_error_code(decode_error::UNSUPPORTED_ENCODING);
}
} else if (dtype_len == 8) {
if (s->dtype_len_in == 4) {
if (s->output_cvt.dtype_len_in == 4) {
// Reading INT32 TIME_MILLIS into 64-bit DURATION_MILLISECONDS
// TIME_MILLIS is the only duration type stored as int32:
// https://github.com/apache/parquet-format/blob/master/LogicalTypes.md#deprecated-time-convertedtype
gpuOutputByteStreamSplit<int32_t>(dst, src, num_values);
// zero out most significant bytes
memset(dst + 4, 0, 4);
} else if (s->ts_scale) {
} else if (s->output_cvt.ts_scale) {
gpuOutputSplitInt64Timestamp(
reinterpret_cast<int64_t*>(dst), src, num_values, s->ts_scale);
reinterpret_cast<int64_t*>(dst), src, num_values, s->output_cvt.ts_scale);
} else {
gpuOutputByteStreamSplit<int64_t>(dst, src, num_values);
}
Expand All @@ -227,8 +227,11 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
auto const& ni = s->nesting_info[leaf_level_index];
if (ni.valid_map != nullptr) {
int const num_values = ni.valid_map_offset - init_valid_map_offset;
zero_fill_null_positions_shared<decode_block_size>(
s, s->dtype_len, init_valid_map_offset, num_values, static_cast<int>(block.thread_rank()));
zero_fill_null_positions_shared<decode_block_size>(s,
s->output_cvt.dtype_len,
init_valid_map_offset,
num_values,
static_cast<int>(block.thread_rank()));
}
}

Expand Down Expand Up @@ -399,7 +402,7 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
// nesting level that is storing actual leaf values
int const leaf_level_index = s->setup.col.max_nesting_depth - 1;

uint32_t const dtype_len = s->dtype_len;
uint32_t const dtype_len = s->output_cvt.dtype_len;
void* dst =
nesting_info_base[leaf_level_index].data_out + static_cast<size_t>(dst_pos) * dtype_len;

Expand Down Expand Up @@ -428,9 +431,9 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
read_fixed_width_value_fast(s, sb, val_src_pos, static_cast<uint2*>(dst));
break;
default:
if (s->dtype_len_in <= sizeof(int32_t)) {
if (s->output_cvt.dtype_len_in <= sizeof(int32_t)) {
read_fixed_width_byte_array_as_int(s, sb, val_src_pos, static_cast<int32_t*>(dst));
} else if (s->dtype_len_in <= sizeof(int64_t)) {
} else if (s->output_cvt.dtype_len_in <= sizeof(int64_t)) {
read_fixed_width_byte_array_as_int(s, sb, val_src_pos, static_cast<int64_t*>(dst));
} else {
read_fixed_width_byte_array_as_int(
Expand All @@ -443,15 +446,15 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
} else if (dtype == Type::INT96) {
read_int96_timestamp(s, sb, val_src_pos, static_cast<int64_t*>(dst));
} else if (dtype_len == 8) {
if (s->dtype_len_in == 4) {
if (s->output_cvt.dtype_len_in == 4) {
// Reading INT32 TIME_MILLIS into 64-bit DURATION_MILLISECONDS
// TIME_MILLIS is the only duration type stored as int32:
// https://github.com/apache/parquet-format/blob/master/LogicalTypes.md#deprecated-time-convertedtype
auto const dst_ptr = static_cast<uint32_t*>(dst);
read_fixed_width_value_fast(s, sb, val_src_pos, dst_ptr);
// zero out most significant bytes
cuda::std::memset(dst_ptr + 1, 0, sizeof(int32_t));
} else if (s->ts_scale) {
} else if (s->output_cvt.ts_scale) {
read_int64_timestamp(s, sb, val_src_pos, static_cast<int64_t*>(dst));
} else {
read_fixed_width_value_fast(s, sb, val_src_pos, static_cast<uint2*>(dst));
Expand All @@ -475,8 +478,11 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
auto const& ni = s->nesting_info[s->setup.col.max_nesting_depth - 1];
if (ni.valid_map != nullptr) {
int const num_values = ni.valid_map_offset - init_valid_map_offset;
zero_fill_null_positions_shared<decode_block_size>(
s, s->dtype_len, init_valid_map_offset, num_values, static_cast<int>(block.thread_rank()));
zero_fill_null_positions_shared<decode_block_size>(s,
s->output_cvt.dtype_len,
init_valid_map_offset,
num_values,
static_cast<int>(block.thread_rank()));
}
}

Expand Down
12 changes: 6 additions & 6 deletions cpp/src/io/parquet/page_data.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -133,7 +133,7 @@ inline __device__ void read_int96_timestamp(page_state_s* s,
dict_pos = src_pos;
src8 = s->stream.data_start;
}
dict_pos *= (uint32_t)s->dtype_len_in;
dict_pos *= (uint32_t)s->output_cvt.dtype_len_in;
ofs = 3 & reinterpret_cast<size_t>(src8);
src8 -= ofs; // align to 32-bit boundary
ofs <<= 3; // bytes -> bits
Expand Down Expand Up @@ -208,7 +208,7 @@ inline __device__ void read_int64_timestamp(page_state_s* s,
dict_pos = src_pos;
src8 = s->stream.data_start;
}
dict_pos *= (uint32_t)s->dtype_len_in;
dict_pos *= (uint32_t)s->output_cvt.dtype_len_in;
ofs = 3 & reinterpret_cast<size_t>(src8);
src8 -= ofs; // align to 32-bit boundary
ofs <<= 3; // bytes -> bits
Expand All @@ -226,7 +226,7 @@ inline __device__ void read_int64_timestamp(page_state_s* s,
val <<= 32;
val |= v.x;
// Output to desired clock rate
ts = apply_ts_scale(val, s->ts_scale);
ts = apply_ts_scale(val, s->output_cvt.ts_scale);
} else {
ts = 0;
}
Expand Down Expand Up @@ -269,7 +269,7 @@ __device__ void read_fixed_width_byte_array_as_int(page_state_s* s,
int src_pos,
T* dst)
{
uint32_t const dtype_len_in = s->dtype_len_in;
uint32_t const dtype_len_in = s->output_cvt.dtype_len_in;
uint8_t const* data = s->stream.dict_base ? s->stream.dict_base : s->stream.data_start;
uint32_t const pos =
(s->stream.dict_base
Expand Down Expand Up @@ -321,7 +321,7 @@ inline __device__ void read_fixed_width_value_fast(page_state_s* s,
dict_pos = src_pos;
dict = s->stream.data_start;
}
dict_pos *= (uint32_t)s->dtype_len_in;
dict_pos *= (uint32_t)s->output_cvt.dtype_len_in;
gpuStoreOutput(dst, dict, dict_pos, dict_size);
}

Expand Down Expand Up @@ -352,7 +352,7 @@ inline __device__ void read_nbyte_fixed_width_value(
dict_pos = src_pos;
dict = s->stream.data_start;
}
dict_pos *= (uint32_t)s->dtype_len_in;
dict_pos *= (uint32_t)s->output_cvt.dtype_len_in;
if (len & 3) {
// Generic slow path
for (unsigned int i = 0; i < len; i++) {
Expand Down
Loading
Loading