From 6d369ec2a6543477d76f3856828971e429783c82 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 10 Aug 2026 11:43:46 -0400 Subject: [PATCH 1/3] Unbind cudf::size_type from offsets used by list columns --- cpp/include/cudf/column/column_factories.hpp | 5 +++- .../cudf/detail/sizes_to_offsets_iterator.cuh | 23 +++++++++------ cpp/include/cudf/lists/detail/gather.cuh | 4 +-- cpp/include/cudf/lists/detail/scatter.cuh | 6 ++-- cpp/include/cudf/lists/list_device_view.cuh | 4 +-- .../cudf/lists/lists_column_device_view.cuh | 2 +- cpp/include/cudf/lists/lists_column_view.hpp | 6 ++-- .../cudf/tdigest/tdigest_column_view.hpp | 2 +- cpp/include/cudf_test/column_wrapper.hpp | 17 ++++++----- cpp/src/groupby/groupby.cu | 2 +- cpp/src/groupby/sort/group_collect.cu | 6 ++-- cpp/src/groupby/sort/group_merge_lists.cu | 6 ++-- cpp/src/hash/md5_hash.cu | 4 +-- cpp/src/io/json/host_tree_algorithms.cu | 2 +- cpp/src/io/json/nested_json.hpp | 8 +++--- cpp/src/io/json/parser_features.cpp | 4 +-- cpp/src/io/orc/stripe_enc.cu | 4 +-- cpp/src/io/orc/writer_impl.cu | 8 +++--- cpp/src/io/statistics/statistics.cuh | 2 +- cpp/src/io/utilities/column_buffer.cpp | 6 ++-- cpp/src/io/utilities/column_buffer_strings.cu | 4 +-- .../combine/concatenate_list_elements.cu | 13 ++++----- cpp/src/lists/combine/concatenate_rows.cu | 14 +++++----- cpp/src/lists/copying/concatenate.cu | 10 +++---- cpp/src/lists/copying/copying.cu | 8 +++--- cpp/src/lists/copying/gather.cu | 10 +++---- cpp/src/lists/copying/scatter_helper.cu | 24 ++++++++-------- cpp/src/lists/dremel.cu | 20 ++++++------- cpp/src/lists/extract.cu | 8 +++--- cpp/src/lists/interleave_columns.cu | 20 ++++++------- cpp/src/lists/lists_column_factories.cu | 8 +++--- cpp/src/lists/lists_column_view.cu | 8 +++--- cpp/src/lists/reverse.cu | 4 +-- cpp/src/lists/segmented_sort.cu | 4 +-- cpp/src/lists/sequences.cu | 4 +-- .../stream_compaction/apply_boolean_mask.cu | 8 +++--- cpp/src/lists/utilities.cu | 15 ++++------ .../rolling/detail/rolling_collect_list.cu | 9 +++--- .../rolling/detail/rolling_collect_list.cuh | 4 +-- cpp/src/rolling/detail/rolling_operators.cuh | 4 +-- cpp/src/strings/convert/convert_lists.cu | 4 +-- cpp/src/strings/repeat_strings.cu | 8 +++--- cpp/src/transform/row_bit_count.cu | 9 +++--- cpp/tests/groupby/collect_list_tests.cpp | 28 ++++++------------- java/src/main/native/src/ColumnViewJni.cpp | 6 ++-- java/src/main/native/src/ColumnViewJni.cu | 4 +-- 46 files changed, 184 insertions(+), 195 deletions(-) diff --git a/cpp/include/cudf/column/column_factories.hpp b/cpp/include/cudf/column/column_factories.hpp index 3423c89d0cee..674aefcd4256 100644 --- a/cpp/include/cudf/column/column_factories.hpp +++ b/cpp/include/cudf/column/column_factories.hpp @@ -484,7 +484,10 @@ std::unique_ptr make_strings_column(size_type num_strings, * @endcode * * @param num_rows The number of lists the column represents. - * @param offsets_column The column of offset values for this column. Each value should + * @param offsets_column The column of offset values for this column. Must be of type + * `type_id::INT32` -- per the Arrow columnar format, a LIST column's offsets are always 32-bit + * (a 64-bit offsets variant would be a distinct LARGE_LIST type, which cudf does not have). + * This is deliberately independent of `cudf::size_type`. Each value should * represent the starting offset into the child elements that corresponds to the beginning of the * row, with the first row starting at 0. The length of row N can be determined by subtracting * `offsets[N+1] - offsets[N]`. The total number of offsets should be 1 longer than the diff --git a/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh b/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh index ae79c6b683d2..14fff8857154 100644 --- a/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh +++ b/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -286,8 +286,10 @@ auto sizes_to_offsets(SizesIterator begin, * The return also includes the total number of elements -- the last element value from the * scan. * + * The returned column is always `type_id::INT32` since offsets children are 32-bit. + * * @throw std::overflow_error if the total size of the scan (last element) greater than maximum - * value of `size_type` + * value of `int32_t` * * @tparam InputIterator Used as input to scan to set the offset values * @param begin The beginning of the input sequence @@ -303,26 +305,29 @@ std::pair, size_type> make_offsets_child_column( rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - auto count = static_cast(std::distance(begin, end)); - auto offsets_column = make_numeric_column( - data_type{type_to_id()}, count + 1, mask_state::UNALLOCATED, stream, mr); + // Offsets children of compound (LIST/STRING) columns are 32-bit and so are + // deliberately not tied to `cudf::size_type`. + auto count = static_cast(std::distance(begin, end)); + auto offsets_column = + make_numeric_column(data_type{type_id::INT32}, count + 1, mask_state::UNALLOCATED, stream, mr); auto offsets_view = offsets_column->mutable_view(); - auto d_offsets = offsets_view.template data(); + auto d_offsets = offsets_view.template data(); // The number of offsets is count+1 so to build the offsets from the sizes // using exclusive-scan technically requires count+1 input values even though // the final input value is never used. // The input iterator is wrapped here to allow the last value to be safely read. auto map_fn = - cuda::proclaim_return_type([begin, count] __device__(size_type idx) -> size_type { - return idx < count ? static_cast(begin[idx]) : size_type{0}; + cuda::proclaim_return_type([begin, count] __device__(size_type idx) -> int32_t { + return idx < count ? static_cast(begin[idx]) : int32_t{0}; }); auto input_itr = cudf::detail::make_counting_transform_iterator(0, map_fn); // Use the sizes-to-offsets iterator to compute the total number of elements auto const total_elements = sizes_to_offsets(input_itr, input_itr + count + 1, d_offsets, 0, stream); + // the offsets are 32-bit so the total must fit in an int32_t CUDF_EXPECTS( - total_elements <= static_cast(std::numeric_limits::max()), + total_elements <= static_cast(std::numeric_limits::max()), "Size of output exceeds the column size limit", std::overflow_error); diff --git a/cpp/include/cudf/lists/detail/gather.cuh b/cpp/include/cudf/lists/detail/gather.cuh index 59ffba0c8367..274888b60019 100644 --- a/cpp/include/cudf/lists/detail/gather.cuh +++ b/cpp/include/cudf/lists/detail/gather.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -100,7 +100,7 @@ gather_data make_gather_data(cudf::lists_column_view const& source_column, // handle sliced columns size_type const shift = source_column.offset() > 0 - ? cudf::detail::get_value(source_column.offsets(), source_column.offset(), stream) + ? cudf::detail::get_value(source_column.offsets(), source_column.offset(), stream) : 0; // generate the base offsets diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index 8c7a3a987dd3..c84c8f67bcfc 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -231,10 +231,10 @@ std::unique_ptr scatter(scalar const& slr, ? cudf::create_null_mask(1, mask_state::UNALLOCATED, stream, mr) : cudf::create_null_mask(1, mask_state::ALL_NULL, stream, mr); auto offset_column = - make_numeric_column(data_type{type_to_id()}, 2, mask_state::UNALLOCATED, stream, mr); + make_numeric_column(data_type{type_id::INT32}, 2, mask_state::UNALLOCATED, stream, mr); thrust::sequence(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - offset_column->mutable_view().begin(), - offset_column->mutable_view().end(), + offset_column->mutable_view().begin(), + offset_column->mutable_view().end(), 0, lv->view().size()); auto wrapped = column_view(data_type{type_id::LIST}, diff --git a/cpp/include/cudf/lists/list_device_view.cuh b/cpp/include/cudf/lists/list_device_view.cuh index dc421d24fa73..894fbbf56d74 100644 --- a/cpp/include/cudf/lists/list_device_view.cuh +++ b/cpp/include/cudf/lists/list_device_view.cuh @@ -40,10 +40,10 @@ class list_device_view { cudf_assert(row_index >= 0 && row_index < lists_column.size() && row_index < offsets.size() && "row_index out of bounds"); - begin_offset = offsets.element(row_index + lists_column.offset()); + begin_offset = offsets.element(row_index + lists_column.offset()); cudf_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() && "begin_offset out of bounds."); - _size = offsets.element(row_index + 1 + lists_column.offset()) - begin_offset; + _size = offsets.element(row_index + 1 + lists_column.offset()) - begin_offset; } ~list_device_view() = default; diff --git a/cpp/include/cudf/lists/lists_column_device_view.cuh b/cpp/include/cudf/lists/lists_column_device_view.cuh index c95463951b08..83be9cb8efa7 100644 --- a/cpp/include/cudf/lists/lists_column_device_view.cuh +++ b/cpp/include/cudf/lists/lists_column_device_view.cuh @@ -79,7 +79,7 @@ class lists_column_device_view : private column_device_view { */ [[nodiscard]] __device__ inline size_type offset_at(size_type idx) const { - return offsets().size() > 0 ? offsets().element(offset() + idx) : 0; + return offsets().size() > 0 ? offsets().element(offset() + idx) : 0; } /** diff --git a/cpp/include/cudf/lists/lists_column_view.hpp b/cpp/include/cudf/lists/lists_column_view.hpp index 0e29ecfab42d..7147462449c5 100644 --- a/cpp/include/cudf/lists/lists_column_view.hpp +++ b/cpp/include/cudf/lists/lists_column_view.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -62,7 +62,7 @@ class lists_column_view : private column_view { using column_view::null_mask; using column_view::offset; using column_view::size; - using offset_iterator = size_type const*; ///< Iterator type for offsets + using offset_iterator = int32_t const*; ///< Iterator type for offsets /** * @brief Returns the parent column. @@ -108,7 +108,7 @@ class lists_column_view : private column_view { */ [[nodiscard]] offset_iterator offsets_begin() const noexcept { - return offsets().begin() + offset(); + return offsets().begin() + offset(); } /** diff --git a/cpp/include/cudf/tdigest/tdigest_column_view.hpp b/cpp/include/cudf/tdigest/tdigest_column_view.hpp index 44848ad3b66b..37458a0c2b79 100644 --- a/cpp/include/cudf/tdigest/tdigest_column_view.hpp +++ b/cpp/include/cudf/tdigest/tdigest_column_view.hpp @@ -66,7 +66,7 @@ class tdigest_column_view : private column_view { tdigest_column_view& operator=(tdigest_column_view&&) = default; using column_view::size; - using offset_iterator = size_type const*; ///< Iterator over offsets + using offset_iterator = int32_t const*; ///< Iterator over offsets // mean and weight column indices within tdigest inner struct columns static constexpr size_type mean_column_index{0}; ///< Mean column index diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp index f0a7a64e3d29..589f18b82ddb 100644 --- a/cpp/include/cudf_test/column_wrapper.hpp +++ b/cpp/include/cudf_test/column_wrapper.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -301,16 +301,15 @@ template auto make_chars_and_offsets(StringsIterator begin, StringsIterator end, ValidityIterator v) { std::vector chars{}; - std::vector offsets(1, 0); + std::vector offsets(1, 0); for (auto str = begin; str < end; ++str) { std::string tmp = (*v++) ? std::string(*str) : std::string{}; chars.insert(chars.end(), std::cbegin(tmp), std::cend(tmp)); auto const last_offset = static_cast(offsets.back()); auto const next_offset = last_offset + tmp.length(); - CUDF_EXPECTS( - next_offset < static_cast(std::numeric_limits::max()), - "Cannot use strings_column_wrapper to build a large strings column"); - offsets.push_back(static_cast(next_offset)); + CUDF_EXPECTS(next_offset < static_cast(std::numeric_limits::max()), + "Cannot use strings_column_wrapper to build a large strings column"); + offsets.push_back(static_cast(next_offset)); } return std::pair(std::move(chars), std::move(offsets)); }; @@ -1544,7 +1543,7 @@ class lists_column_wrapper : public detail::column_wrapper { */ static lists_column_wrapper make_one_empty_row_column(bool valid = true) { - cudf::test::fixed_width_column_wrapper offsets{0, 0}; + cudf::test::fixed_width_column_wrapper offsets{0, 0}; cudf::test::fixed_width_column_wrapper values{}; return lists_column_wrapper( 1, @@ -1627,7 +1626,7 @@ class lists_column_wrapper : public detail::column_wrapper { // add the final offset offsetv.push_back(count); auto offsets = - cudf::test::fixed_width_column_wrapper(offsetv.begin(), offsetv.end()).release(); + cudf::test::fixed_width_column_wrapper(offsetv.begin(), offsetv.end()).release(); // concatenate them together, skipping children that are null. std::vector children; @@ -1673,7 +1672,7 @@ class lists_column_wrapper : public detail::column_wrapper { offsetv.push_back(c->size()); } auto offsets = - cudf::test::fixed_width_column_wrapper(offsetv.begin(), offsetv.end()).release(); + cudf::test::fixed_width_column_wrapper(offsetv.begin(), offsetv.end()).release(); // construct the list column. mark this as a root root = true; diff --git a/cpp/src/groupby/groupby.cu b/cpp/src/groupby/groupby.cu index cf61064724f8..de7079259904 100644 --- a/cpp/src/groupby/groupby.cu +++ b/cpp/src/groupby/groupby.cu @@ -104,7 +104,7 @@ struct empty_column_constructor { if constexpr (k == aggregation::Kind::HISTOGRAM) { return make_lists_column(0, - make_empty_column(type_to_id()), + make_empty_column(type_id::INT32), cudf::reduction::detail::make_empty_histogram_like(values), 0, {}); diff --git a/cpp/src/groupby/sort/group_collect.cu b/cpp/src/groupby/sort/group_collect.cu index 992cca53a456..9475508a1550 100644 --- a/cpp/src/groupby/sort/group_collect.cu +++ b/cpp/src/groupby/sort/group_collect.cu @@ -63,7 +63,7 @@ std::pair, std::unique_ptr> purge_null_entries( cuda::counting_iterator{0}, cuda::counting_iterator{num_groups}, null_purged_sizes.begin(), - [d_offsets = offsets.template begin(), not_null_pred] __device__(auto i) { + [d_offsets = offsets.template begin(), not_null_pred] __device__(auto i) { return thrust::count_if(thrust::seq, cuda::counting_iterator{d_offsets[i]}, cuda::counting_iterator{d_offsets[i + 1]}, @@ -86,12 +86,12 @@ std::unique_ptr group_collect(column_view const& values, auto [child_column, offsets_column] = [null_handling, num_groups, &values, &group_offsets, stream, mr] { auto offsets_column = make_numeric_column( - data_type(type_to_id()), num_groups + 1, mask_state::UNALLOCATED, stream, mr); + data_type(type_id::INT32), num_groups + 1, mask_state::UNALLOCATED, stream, mr); thrust::copy(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), group_offsets.begin(), group_offsets.end(), - offsets_column->mutable_view().template begin()); + offsets_column->mutable_view().template begin()); // If column of grouped values contains null elements, and null_policy == EXCLUDE, // those elements must be filtered out, and offsets recomputed. diff --git a/cpp/src/groupby/sort/group_merge_lists.cu b/cpp/src/groupby/sort/group_merge_lists.cu index 7dd04732f22d..3cea3b6a5de2 100644 --- a/cpp/src/groupby/sort/group_merge_lists.cu +++ b/cpp/src/groupby/sort/group_merge_lists.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -28,7 +28,7 @@ std::unique_ptr group_merge_lists(column_view const& values, "Input to `group_merge_lists` must be a non-nullable lists column."); auto offsets_column = make_numeric_column( - data_type(type_to_id()), num_groups + 1, mask_state::UNALLOCATED, stream, mr); + data_type(type_id::INT32), num_groups + 1, mask_state::UNALLOCATED, stream, mr); // Generate offsets of the output lists column by gathering from the provided group offsets and // the input list offsets. @@ -44,7 +44,7 @@ std::unique_ptr group_merge_lists(column_view const& values, group_offsets.begin(), group_offsets.end(), lists_column_view(values).offsets_begin(), - offsets_column->mutable_view().template begin()); + offsets_column->mutable_view().template begin()); // The child column of the output lists column is just copied from the input column. auto child_column = diff --git a/cpp/src/hash/md5_hash.cu b/cpp/src/hash/md5_hash.cu index 0d42eac3870d..2ba82faebd8a 100644 --- a/cpp/src/hash/md5_hash.cu +++ b/cpp/src/hash/md5_hash.cu @@ -323,8 +323,8 @@ std::unique_ptr md5(table_view const& input, if (data_col.type().id() == type_id::LIST) { CUDF_UNREACHABLE("Nested list unsupported"); } - auto const offset_begin = offsets.element(row_index); - auto const offset_end = offsets.element(row_index + 1); + auto const offset_begin = offsets.element(row_index); + auto const offset_end = offsets.element(row_index + 1); cudf::type_dispatcher( data_col.type(), ListHasherDispatcher(&hasher, data_col), offset_begin, offset_end); } else { diff --git a/cpp/src/io/json/host_tree_algorithms.cu b/cpp/src/io/json/host_tree_algorithms.cu index cfa616f6ec45..a85b996ac55f 100644 --- a/cpp/src/io/json/host_tree_algorithms.cu +++ b/cpp/src/io/json/host_tree_algorithms.cu @@ -213,7 +213,7 @@ struct json_column_data { using row_offset_t = json_column::row_offset_t; row_offset_t* string_offsets; row_offset_t* string_lengths; - row_offset_t* child_offsets; + int32_t* child_offsets; bitmask_type* validity; }; diff --git a/cpp/src/io/json/nested_json.hpp b/cpp/src/io/json/nested_json.hpp index b7845c22ce64..de9368a6bb74 100644 --- a/cpp/src/io/json/nested_json.hpp +++ b/cpp/src/io/json/nested_json.hpp @@ -72,8 +72,8 @@ struct json_column { std::vector string_offsets; std::vector string_lengths; - // Row offsets - std::vector child_offsets; + // Row offsets (LIST offsets child data; always 32-bit) + std::vector child_offsets; // Validity bitmap std::vector validity; @@ -148,8 +148,8 @@ struct device_json_column { rmm::device_uvector string_offsets; rmm::device_uvector string_lengths; - // Row offsets - rmm::device_uvector child_offsets; + // Row offsets (LIST offsets child data; always 32-bit) + rmm::device_uvector child_offsets; // Validity bitmap rmm::device_buffer validity; diff --git a/cpp/src/io/json/parser_features.cpp b/cpp/src/io/json/parser_features.cpp index 68869579b5da..1ee13c1f747d 100644 --- a/cpp/src/io/json/parser_features.cpp +++ b/cpp/src/io/json/parser_features.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -75,7 +75,7 @@ struct empty_column_functor { auto const& child_name = schema.child_types.begin()->first; std::unique_ptr child = cudf::type_dispatcher( schema.child_types.at(child_name).type, *this, schema.child_types.at(child_name)); - auto offsets = make_empty_column(data_type(type_to_id())); + auto offsets = make_empty_column(data_type(type_id::INT32)); std::vector> child_columns; child_columns.push_back(std::move(offsets)); child_columns.push_back(std::move(child)); diff --git a/cpp/src/io/orc/stripe_enc.cu b/cpp/src/io/orc/stripe_enc.cu index d76c0dcc57ee..e03c331981de 100644 --- a/cpp/src/io/orc/stripe_enc.cu +++ b/cpp/src/io/orc/stripe_enc.cu @@ -845,8 +845,8 @@ CUDF_KERNEL void __launch_bounds__(block_size) case MAP: { auto const& offsets = column.child(lists_column_view::offsets_column_index); // Compute list length from the offsets - s->lengths.u32[nz_idx] = offsets.element(row + 1 + column.offset()) - - offsets.element(row + column.offset()); + s->lengths.u32[nz_idx] = offsets.element(row + 1 + column.offset()) - + offsets.element(row + column.offset()); } break; default: break; } diff --git a/cpp/src/io/orc/writer_impl.cu b/cpp/src/io/orc/writer_impl.cu index 62ff8f165c3e..60f15af0446a 100644 --- a/cpp/src/io/orc/writer_impl.cu +++ b/cpp/src/io/orc/writer_impl.cu @@ -1703,8 +1703,8 @@ void pushdown_lists_null_mask(orc_column_view const& col, auto const is_row_valid = d_col.is_valid(idx) and bit_value_or(parent_pd_mask, idx, true); if (not is_row_valid) { auto offsets = d_col.child(lists_column_view::offsets_column_index); - auto const child_rows_begin = offsets.element(idx + d_col.offset()); - auto const child_rows_end = offsets.element(idx + 1 + d_col.offset()); + auto const child_rows_begin = offsets.element(idx + d_col.offset()); + auto const child_rows_end = offsets.element(idx + 1 + d_col.offset()); for (auto child_row = child_rows_begin; child_row < child_rows_end; ++child_row) clear_bit(out_mask.data(), child_row); } @@ -1953,9 +1953,9 @@ hostdevice_2dvector calculate_rowgroup_bounds(orc_table_view cons auto offsets = parent_col.child(lists_column_view::offsets_column_index); auto const rows_begin = - offsets.element(parent_rg.begin + parent_col.offset()) - col.offset(); + offsets.element(parent_rg.begin + parent_col.offset()) - col.offset(); auto const rows_end = - offsets.element(parent_rg.end + parent_col.offset()) - col.offset(); + offsets.element(parent_rg.end + parent_col.offset()) - col.offset(); return rowgroup_rows{rows_begin, rows_end}; } diff --git a/cpp/src/io/statistics/statistics.cuh b/cpp/src/io/statistics/statistics.cuh index 1897cf272a2c..3526a8824ead 100644 --- a/cpp/src/io/statistics/statistics.cuh +++ b/cpp/src/io/statistics/statistics.cuh @@ -118,7 +118,7 @@ __device__ T get_element(column_device_view const& col, uint32_t row) { using et = typename T::element_type; size_type const index = row + col.offset(); // account for this view's _offset - auto const* d_offsets = col.child(lists_column_view::offsets_column_index).data(); + auto const* d_offsets = col.child(lists_column_view::offsets_column_index).data(); auto const* d_data = col.child(lists_column_view::child_column_index).data(); auto const offset = d_offsets[index]; return T(d_data + offset, d_offsets[index + 1] - offset); diff --git a/cpp/src/io/utilities/column_buffer.cpp b/cpp/src/io/utilities/column_buffer.cpp index c1af6be6acd4..ea5facad4a34 100644 --- a/cpp/src/io/utilities/column_buffer.cpp +++ b/cpp/src/io/utilities/column_buffer.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -52,7 +52,7 @@ void cudf::io::detail::inline_column_buffer::allocate_strings_data(bool memset_d { CUDF_EXPECTS(type.id() == type_id::STRING, "allocate_strings_data called for non-string column"); // size + 1 for final offset. _string_data will be initialized later. - _data = create_data(data_type{type_to_id()}, size + 1, memset_data, stream, _mr); + _data = create_data(data_type{type_id::INT32}, size + 1, memset_data, stream, _mr); } void cudf::io::detail::inline_column_buffer::create_string_data(size_t num_bytes, @@ -107,7 +107,7 @@ void column_buffer_base::create_with_mask(size_type _size, // list columns store a buffer of int32's as offsets to represent // their individual rows case type_id::LIST: - _data = create_data(data_type{type_to_id()}, size, memset_data, stream, _mr); + _data = create_data(data_type{type_id::INT32}, size, memset_data, stream, _mr); break; // struct columns store no data themselves. just validity and children. diff --git a/cpp/src/io/utilities/column_buffer_strings.cu b/cpp/src/io/utilities/column_buffer_strings.cu index eba7c45c0de3..1d50fb4d5d69 100644 --- a/cpp/src/io/utilities/column_buffer_strings.cu +++ b/cpp/src/io/utilities/column_buffer_strings.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -33,7 +33,7 @@ std::unique_ptr cudf::io::detail::inline_column_buffer::make_string_colu } else { // no need for copies, just transfer ownership of the data_buffers to the columns auto offsets_col = std::make_unique( - data_type{type_to_id()}, size + 1, std::move(_data), rmm::device_buffer{}, 0); + data_type{type_id::INT32}, size + 1, std::move(_data), rmm::device_buffer{}, 0); return make_strings_column( size, std::move(offsets_col), std::move(_string_data), null_count(), std::move(_null_mask)); } diff --git a/cpp/src/lists/combine/concatenate_list_elements.cu b/cpp/src/lists/combine/concatenate_list_elements.cu index a473f4a890d3..ebcc183437e4 100644 --- a/cpp/src/lists/combine/concatenate_list_elements.cu +++ b/cpp/src/lists/combine/concatenate_list_elements.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -47,9 +47,9 @@ std::unique_ptr concatenate_lists_ignore_null(column_view const& input, auto const num_rows = input.size(); auto out_offsets = make_numeric_column( - data_type{type_to_id()}, num_rows + 1, mask_state::UNALLOCATED, stream, mr); + data_type{type_id::INT32}, num_rows + 1, mask_state::UNALLOCATED, stream, mr); - auto const d_out_offsets = out_offsets->mutable_view().template begin(); + auto const d_out_offsets = out_offsets->mutable_view().template begin(); auto const d_row_offsets = lists_column_view(input).offsets_begin(); auto const d_list_offsets = lists_column_view(lists_column_view(input).child()).offsets_begin(); @@ -179,7 +179,7 @@ std::unique_ptr gather_list_entries(column_view const& input, d_list_offsets, d_indices = gather_map.begin(), d_out_list_offsets = - output_list_offsets.template begin()] __device__(size_type const idx) { + output_list_offsets.template begin()] __device__(size_type const idx) { // The output row has been identified as a null/empty list during list size computation. if (d_out_list_offsets[idx + 1] == d_out_list_offsets[idx]) { return; } @@ -207,9 +207,8 @@ std::unique_ptr concatenate_lists_nullifying_rows(column_view const& inp auto [list_offsets, list_validities] = generate_list_offsets_and_validities(input, stream, mr); auto const offsets_view = list_offsets->view(); - auto const num_rows = input.size(); - auto const num_output_entries = - cudf::detail::get_value(offsets_view, num_rows, stream); + auto const num_rows = input.size(); + auto const num_output_entries = cudf::detail::get_value(offsets_view, num_rows, stream); auto list_entries = gather_list_entries(input, offsets_view, num_rows, num_output_entries, stream, mr); diff --git a/cpp/src/lists/combine/concatenate_rows.cu b/cpp/src/lists/combine/concatenate_rows.cu index 2964276a3fbc..7fa0bf509846 100644 --- a/cpp/src/lists/combine/concatenate_rows.cu +++ b/cpp/src/lists/combine/concatenate_rows.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -69,7 +69,7 @@ generate_regrouped_offsets_and_null_mask(table_device_view const& input, { // outgoing offsets. auto offsets = cudf::make_fixed_width_column( - data_type{type_to_id()}, input.num_rows() + 1, mask_state::UNALLOCATED, stream, mr); + data_type{type_id::INT32}, input.num_rows() + 1, mask_state::UNALLOCATED, stream, mr); auto keys = thrust::make_transform_iterator(cuda::counting_iterator{0}, @@ -96,7 +96,7 @@ generate_regrouped_offsets_and_null_mask(table_device_view const& input, } } auto offsets = - input.column(col_index).child(lists_column_view::offsets_column_index).data() + + input.column(col_index).child(lists_column_view::offsets_column_index).data() + input.column(col_index).offset(); return offsets[row_index + 1] - offsets[row_index]; })); @@ -105,15 +105,15 @@ generate_regrouped_offsets_and_null_mask(table_device_view const& input, keys + (input.num_rows() * input.num_columns()), values, cuda::make_discard_iterator(), - offsets->mutable_view().begin(), + offsets->mutable_view().begin(), cuda::std::plus(), stream); // convert to offsets thrust::exclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - offsets->view().begin(), - offsets->view().begin() + input.num_rows() + 1, - offsets->mutable_view().begin(), + offsets->view().begin(), + offsets->view().begin() + input.num_rows() + 1, + offsets->mutable_view().begin(), 0); // generate appropriate null mask diff --git a/cpp/src/lists/copying/concatenate.cu b/cpp/src/lists/copying/concatenate.cu index 8c4e2826672e..3324f18ec3c2 100644 --- a/cpp/src/lists/copying/concatenate.cu +++ b/cpp/src/lists/copying/concatenate.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -63,13 +63,13 @@ std::unique_ptr merge_offsets(host_span columns // handle sliced columns int const local_shift = shift - - (c.offset() > 0 ? cudf::detail::get_value(c.offsets(), c.offset(), stream) : 0); + (c.offset() > 0 ? cudf::detail::get_value(c.offsets(), c.offset(), stream) : 0); column_device_view offsets(c.offsets(), nullptr, nullptr); thrust::transform( rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - offsets.begin() + c.offset(), - offsets.begin() + c.offset() + c.size() + 1, - d_merged_offsets.begin() + count, + offsets.begin() + c.offset(), + offsets.begin() + c.offset() + c.size() + 1, + d_merged_offsets.begin() + count, [local_shift] __device__(size_type offset) { return offset + local_shift; }); shift += c.get_sliced_child(stream).size(); diff --git a/cpp/src/lists/copying/copying.cu b/cpp/src/lists/copying/copying.cu index 37a7054c2fa5..a4bcf7688bdd 100644 --- a/cpp/src/lists/copying/copying.cu +++ b/cpp/src/lists/copying/copying.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include @@ -40,9 +40,9 @@ std::unique_ptr copy_slice(lists_column_view const& lists, end += lists.offset(); // Offsets at the beginning and end of the slice: - auto offsets_data = lists.offsets().data(); - auto start_offset = cudf::detail::get_value(lists.offsets(), start, stream); - auto end_offset = cudf::detail::get_value(lists.offsets(), end, stream); + auto offsets_data = lists.offsets().data(); + auto start_offset = cudf::detail::get_value(lists.offsets(), start, stream); + auto end_offset = cudf::detail::get_value(lists.offsets(), end, stream); rmm::device_uvector out_offsets(offsets_count, stream); diff --git a/cpp/src/lists/copying/gather.cu b/cpp/src/lists/copying/gather.cu index 4255cf7478fa..0601b2138678 100644 --- a/cpp/src/lists/copying/gather.cu +++ b/cpp/src/lists/copying/gather.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -50,20 +50,20 @@ struct list_gatherer { using result_type = size_type; size_t offset_count; - size_type const* base_offsets; - size_type const* offsets; + int32_t const* base_offsets; + int32_t const* offsets; list_gatherer(gather_data const& gd) : offset_count{gd.base_offsets.size()}, base_offsets{gd.base_offsets.data()}, - offsets{gd.offsets->mutable_view().data()} + offsets{gd.offsets->mutable_view().data()} { } __device__ result_type operator()(argument_type index) { // the "upper bound" of the span for a given offset is always offsets+1; - size_type const* upper_bound_start = offsets + 1; + int32_t const* upper_bound_start = offsets + 1; // "step 1" from above auto const bound = thrust::upper_bound(thrust::seq, upper_bound_start, upper_bound_start + offset_count, index); diff --git a/cpp/src/lists/copying/scatter_helper.cu b/cpp/src/lists/copying/scatter_helper.cu index a6295576b0fa..5c009804eff9 100644 --- a/cpp/src/lists/copying/scatter_helper.cu +++ b/cpp/src/lists/copying/scatter_helper.cu @@ -47,7 +47,7 @@ std::pair construct_child_nullmask( rmm::device_async_resource_ref mr) { auto is_valid_predicate = [d_list_vector = parent_list_vector.begin(), - d_offsets = parent_list_offsets.template data(), + d_offsets = parent_list_offsets.template data(), d_offsets_size = parent_list_offsets.size(), source_lists, target_lists] __device__(auto const& i) { @@ -160,7 +160,7 @@ struct list_child_constructor { auto target_lists = cudf::lists_column_device_view(*target_column_device_view); auto const num_child_rows{ - cudf::detail::get_value(list_offsets, list_offsets.size() - 1, stream)}; + cudf::detail::get_value(list_offsets, list_offsets.size() - 1, stream)}; auto child_null_mask = source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() @@ -180,7 +180,7 @@ struct list_child_constructor { cuda::counting_iterator{0}, cuda::counting_iterator{child_column->size()}, child_column->mutable_view().begin(), - cuda::proclaim_return_type([offset_begin = list_offsets.begin(), + cuda::proclaim_return_type([offset_begin = list_offsets.begin(), offset_size = list_offsets.size(), d_list_vector = list_vector.begin(), source_lists, @@ -219,7 +219,7 @@ struct list_child_constructor { auto target_lists = cudf::lists_column_device_view(*target_column_device_view); auto const num_child_rows{ - cudf::detail::get_value(list_offsets, list_offsets.size() - 1, stream)}; + cudf::detail::get_value(list_offsets, list_offsets.size() - 1, stream)}; if (num_child_rows == 0) { return make_empty_column(type_id::STRING); } @@ -232,7 +232,7 @@ struct list_child_constructor { cuda::counting_iterator{0}, cuda::counting_iterator{static_cast(string_views.size())}, string_views.begin(), - cuda::proclaim_return_type([offset_begin = list_offsets.begin(), + cuda::proclaim_return_type([offset_begin = list_offsets.begin(), offset_size = list_offsets.size(), d_list_vector = list_vector.begin(), source_lists, @@ -246,7 +246,7 @@ struct list_child_constructor { auto row_index = d_list_vector[list_index].row_index(); auto actual_list_row = d_list_vector[list_index].bind_to_column(source_lists, target_lists); auto lists_column = actual_list_row.get_column(); - auto lists_offsets_ptr = lists_column.offsets().template data(); + auto lists_offsets_ptr = lists_column.offsets().template data(); auto child_strings_column = lists_column.child(); auto strings_offset = lists_offsets_ptr[row_index] + intra_index; @@ -282,7 +282,7 @@ struct list_child_constructor { auto target_lists = cudf::lists_column_device_view(*target_column_device_view); auto const num_child_rows{ - cudf::detail::get_value(list_offsets, list_offsets.size() - 1, stream)}; + cudf::detail::get_value(list_offsets, list_offsets.size() - 1, stream)}; if (num_child_rows == 0) { // make an empty lists column using the input child type @@ -299,8 +299,8 @@ struct list_child_constructor { cuda::counting_iterator{0}, cuda::counting_iterator{static_cast(child_list_views.size())}, child_list_views.begin(), - cuda::proclaim_return_type([offset_begin = list_offsets.begin(), - offset_size = list_offsets.size(), + cuda::proclaim_return_type([offset_begin = list_offsets.begin(), + offset_size = list_offsets.size(), d_list_vector = list_vector.begin(), source_lists, target_lists] __device__(auto index) { @@ -314,10 +314,10 @@ struct list_child_constructor { auto actual_list_row = d_list_vector[list_index].bind_to_column(source_lists, target_lists); auto lists_column = actual_list_row.get_column(); auto child_lists_column = lists_column.child(); - auto lists_offsets_ptr = lists_column.offsets().template data(); + auto lists_offsets_ptr = lists_column.offsets().template data(); auto child_lists_offsets_ptr = child_lists_column.child(lists_column_view::offsets_column_index) - .template data(); + .template data(); auto child_row_index = lists_offsets_ptr[row_index] + intra_index; auto size = child_lists_offsets_ptr[child_row_index + 1] - child_lists_offsets_ptr[child_row_index]; @@ -379,7 +379,7 @@ struct list_child_constructor { auto const target_structs = target_lists_column_view.child(); auto const num_child_rows{ - cudf::detail::get_value(list_offsets, list_offsets.size() - 1, stream)}; + cudf::detail::get_value(list_offsets, list_offsets.size() - 1, stream)}; auto const num_struct_members = std::distance(source_structs.child_begin(), source_structs.child_end()); diff --git a/cpp/src/lists/dremel.cu b/cpp/src/lists/dremel.cu index b9f6cf2e3450..ec477853656c 100644 --- a/cpp/src/lists/dremel.cu +++ b/cpp/src/lists/dremel.cu @@ -82,7 +82,7 @@ dremel_data get_encoding(column_view h_col, auto lcv = lists_column_view(get_list_level(col)); rmm::device_uvector empties_idx(lcv.size(), stream); rmm::device_uvector empties(lcv.size(), stream); - auto d_off = lcv.offsets().data(); + auto d_off = lcv.offsets().data(); auto empties_idx_end = cudf::detail::copy_if( cuda::counting_iterator{start}, @@ -94,7 +94,7 @@ dremel_data get_encoding(column_view h_col, thrust::gather(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), empties_idx.begin(), empties_idx_end, - lcv.offsets().begin(), + lcv.offsets().begin(), empties.begin()); auto empties_size = empties_end - empties.begin(); @@ -120,13 +120,13 @@ dremel_data get_encoding(column_view h_col, } std::unique_ptr empty_list_offset_col; if (has_empty_list_offsets) { - empty_list_offset_col = make_fixed_width_column(data_type(type_to_id()), + empty_list_offset_col = make_fixed_width_column(data_type(type_id::INT32), 1, mask_state::UNALLOCATED, stream, cudf::get_current_device_resource_ref()); CUDF_CUDA_TRY(cudaMemsetAsync( - empty_list_offset_col->mutable_view().head(), 0, sizeof(size_type), stream.value())); + empty_list_offset_col->mutable_view().head(), 0, sizeof(int32_t), stream.value())); std::function normalize_col = [&](column_view const& col) { auto children = [&]() -> std::vector { if (col.type().id() == type_id::LIST) { @@ -238,8 +238,8 @@ dremel_data get_encoding(column_view h_col, // Skip doing the following for any structs we encounter in between. while (curr_col.type().id() == type_id::LIST or curr_col.type().id() == type_id::STRUCT) { if (curr_col.type().id() == type_id::LIST) { - off = curr_col.child(lists_column_view::offsets_column_index).element(off); - end = curr_col.child(lists_column_view::offsets_column_index).element(end); + off = curr_col.child(lists_column_view::offsets_column_index).element(off); + end = curr_col.child(lists_column_view::offsets_column_index).element(end); if (level < level_max) { offset_at_level[level] = off; end_idx_at_level[level] = end; @@ -335,7 +335,7 @@ dremel_data get_encoding(column_view h_col, // Scan to get distance by which each offset value is shifted due to the insertion of empties auto scan_it = cudf::detail::make_counting_transform_iterator( column_offsets[level], - cuda::proclaim_return_type([off = lcv.offsets().data(), + cuda::proclaim_return_type([off = lcv.offsets().data(), size = lcv.offsets().size()] __device__(auto i) -> int { return (i + 1 < size) && (off[i] == off[i + 1]); })); @@ -350,7 +350,7 @@ dremel_data get_encoding(column_view h_col, thrust::for_each_n(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), cuda::counting_iterator{0}, offset_size_at_level, - [off = lcv.offsets().data() + column_offsets[level], + [off = lcv.offsets().data() + column_offsets[level], scan_out = scan_out.data(), new_off = new_offsets.data()] __device__(auto i) { new_off[i] = off[i] - off[0] + scan_out[i]; @@ -426,7 +426,7 @@ dremel_data get_encoding(column_view h_col, // level value fof an empty list auto scan_it = cudf::detail::make_counting_transform_iterator( column_offsets[level], - cuda::proclaim_return_type([off = lcv.offsets().data(), + cuda::proclaim_return_type([off = lcv.offsets().data(), size = lcv.offsets().size()] __device__(auto i) -> int { return (i + 1 < size) && (off[i] == off[i + 1]); })); @@ -441,7 +441,7 @@ dremel_data get_encoding(column_view h_col, thrust::for_each_n(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), cuda::counting_iterator{0}, offset_size_at_level, - [off = lcv.offsets().data() + column_offsets[level], + [off = lcv.offsets().data() + column_offsets[level], scan_out = scan_out.data(), new_off = temp_new_offsets.data(), offset_transformer] __device__(auto i) { diff --git a/cpp/src/lists/extract.cu b/cpp/src/lists/extract.cu index bc9cfa0d3991..7f794a4e1c72 100644 --- a/cpp/src/lists/extract.cu +++ b/cpp/src/lists/extract.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -54,7 +54,7 @@ std::unique_ptr make_index_child(column_view const& indices, // Replace null indices with MAX_SIZE_TYPE, so that gather() returns null for them. auto const null_replaced_iter_begin = cudf::detail::make_null_replacement_iterator(*d_indices, std::numeric_limits::max()); - auto index_child = make_numeric_column(data_type{type_id::INT32}, + auto index_child = make_numeric_column(data_type{type_to_id()}, indices.size(), mask_state::UNALLOCATED, stream, @@ -80,7 +80,7 @@ std::unique_ptr make_index_child(size_type index, rmm::cuda_stream_view stream) { auto index_child = // [index, index, index, ..., index] - make_numeric_column(data_type{type_id::INT32}, + make_numeric_column(data_type{type_to_id()}, num_rows, mask_state::UNALLOCATED, stream, @@ -103,7 +103,7 @@ std::unique_ptr make_index_offsets(size_type num_lists, rmm::cuda_ { return cudf::detail::sequence( num_lists + 1, - cudf::scalar_type_t(0, true, stream, cudf::get_current_device_resource_ref()), + cudf::scalar_type_t(0, true, stream, cudf::get_current_device_resource_ref()), stream, cudf::get_current_device_resource_ref()); } diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 202dc1183418..76c1c115a5dd 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -49,8 +49,8 @@ generate_list_offsets_and_validities(table_view const& input, // The output offsets column. auto list_offsets = make_numeric_column( - data_type{type_to_id()}, num_output_lists + 1, mask_state::UNALLOCATED, stream, mr); - auto const d_offsets = list_offsets->mutable_view().template begin(); + data_type{type_id::INT32}, num_output_lists + 1, mask_state::UNALLOCATED, stream, mr); + auto const d_offsets = list_offsets->mutable_view().template begin(); // The array of int8_t to store validities for list elements. auto validities = rmm::device_uvector(has_null_mask ? num_output_lists : 0, stream); @@ -70,7 +70,7 @@ generate_list_offsets_and_validities(table_view const& input, auto const& lists_col = table_dv.column(col_id); if (has_null_mask) { d_validities[idx] = static_cast(lists_col.is_valid(list_id)); } auto const list_offsets = - lists_col.child(lists_column_view::offsets_column_index).template data() + + lists_col.child(lists_column_view::offsets_column_index).template data() + lists_col.offset(); return list_offsets[list_id + 1] - list_offsets[list_id]; })); @@ -138,7 +138,7 @@ struct compute_string_sizes_and_interleave_lists_fn { table_device_view const table_dv; // Store list offsets of the output lists column. - size_type const* const dst_list_offsets; + int32_t const* const dst_list_offsets; using string_index_pair = cudf::strings::detail::string_index_pair; string_index_pair* indices; // output @@ -154,7 +154,7 @@ struct compute_string_sizes_and_interleave_lists_fn { if (lists_col.is_null(list_id)) { return; } auto const list_offsets = - lists_col.child(lists_column_view::offsets_column_index).template data() + + lists_col.child(lists_column_view::offsets_column_index).template data() + lists_col.offset(); auto const& str_col = lists_col.child(lists_column_view::child_column_index); @@ -192,7 +192,7 @@ struct interleave_list_entries_impl(); + auto const d_list_offsets = output_list_offsets.template begin(); rmm::device_uvector indices(num_output_entries, stream); @@ -238,14 +238,14 @@ struct interleave_list_entries_impl( [num_cols, table_dv = *table_dv_ptr, d_validities = validities.begin(), - d_offsets = output_list_offsets.template begin(), + d_offsets = output_list_offsets.template begin(), d_output = output_dv_ptr->template begin(), data_has_null_mask] __device__(size_type const idx) { auto const col_id = idx % num_cols; auto const list_id = idx / num_cols; auto const& lists_col = table_dv.column(col_id); auto const list_offsets = - lists_col.child(lists_column_view::offsets_column_index).template data() + + lists_col.child(lists_column_view::offsets_column_index).template data() + lists_col.offset(); auto const& data_col = lists_col.child(lists_column_view::child_column_index); @@ -346,7 +346,7 @@ std::unique_ptr interleave_columns(table_view const& input, // specialized for different types. auto const num_output_lists = input.num_rows() * input.num_columns(); auto const num_output_entries = - cudf::detail::get_value(offsets_view, num_output_lists, stream); + cudf::detail::get_value(offsets_view, num_output_lists, stream); auto const data_has_null_mask = std::any_of(std::cbegin(input), std::cend(input), [](auto const& col) { return col.child(lists_column_view::child_column_index).nullable(); diff --git a/cpp/src/lists/lists_column_factories.cu b/cpp/src/lists/lists_column_factories.cu index 51cd76764b36..26306bb32ca9 100644 --- a/cpp/src/lists/lists_column_factories.cu +++ b/cpp/src/lists/lists_column_factories.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -30,7 +30,7 @@ std::unique_ptr make_lists_column_from_scalar(list_scalar const& v if (size == 0) { return make_lists_column( 0, - make_empty_column(type_to_id()), + make_empty_column(type_id::INT32), empty_like(value.view()), 0, cudf::detail::create_null_mask(0, mask_state::UNALLOCATED, stream, mr)); @@ -72,7 +72,7 @@ std::unique_ptr make_lists_column_from_scalar(list_scalar const& v std::unique_ptr make_empty_lists_column(data_type child_type) { - auto offsets = make_empty_column(data_type(type_to_id())); + auto offsets = make_empty_column(data_type(type_id::INT32)); auto child = make_empty_column(child_type); return make_lists_column(0, std::move(offsets), std::move(child), 0, rmm::device_buffer{}); } @@ -84,7 +84,7 @@ std::unique_ptr make_all_nulls_lists_column(size_type size, { auto offsets = [&] { auto offsets_buff = - cudf::detail::make_zeroed_device_uvector_async(size + 1, stream, mr); + cudf::detail::make_zeroed_device_uvector_async(size + 1, stream, mr); return std::make_unique(std::move(offsets_buff), rmm::device_buffer{}, 0); }(); auto child = make_empty_column(child_type); diff --git a/cpp/src/lists/lists_column_view.cu b/cpp/src/lists/lists_column_view.cu index 7edd19d51b6f..51d56c0dd724 100644 --- a/cpp/src/lists/lists_column_view.cu +++ b/cpp/src/lists/lists_column_view.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2022, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -38,16 +38,16 @@ column_view lists_column_view::get_sliced_child(rmm::cuda_stream_view stream) co if (offset() > 0) { // theoretically this function could always do this step and be correct, but get_value<> // actually hits the gpu so it's best to avoid it if possible. - size_type child_offset_start = cudf::detail::get_value(offsets(), offset(), stream); + size_type child_offset_start = cudf::detail::get_value(offsets(), offset(), stream); size_type child_offset_end = - cudf::detail::get_value(offsets(), offset() + size(), stream); + cudf::detail::get_value(offsets(), offset() + size(), stream); return cudf::detail::slice(child(), {child_offset_start, child_offset_end}, stream).front(); } // if I don't have a positive offset, but I am shorter than my offsets() would otherwise indicate, // I need to do a split and return the front. if (size() < offsets().size() - 1) { - size_type child_offset = cudf::detail::get_value(offsets(), size(), stream); + size_type child_offset = cudf::detail::get_value(offsets(), size(), stream); return cudf::detail::slice(child(), {0, child_offset}, stream).front(); } diff --git a/cpp/src/lists/reverse.cu b/cpp/src/lists/reverse.cu index f0c1e5975c74..23228de7cd22 100644 --- a/cpp/src/lists/reverse.cu +++ b/cpp/src/lists/reverse.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -46,7 +46,7 @@ std::unique_ptr reverse(lists_column_view const& input, thrust::for_each_n(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), cuda::counting_iterator{0}, child.size(), - [list_offsets = out_offsets->view().begin(), + [list_offsets = out_offsets->view().begin(), list_indices = labels->view().begin(), gather_map = gather_map.begin()] __device__(auto const idx) { auto const list_idx = list_indices[idx]; diff --git a/cpp/src/lists/segmented_sort.cu b/cpp/src/lists/segmented_sort.cu index a10df52104a2..1e246a062d76 100644 --- a/cpp/src/lists/segmented_sort.cu +++ b/cpp/src/lists/segmented_sort.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -40,7 +40,7 @@ std::unique_ptr build_output_offsets(lists_column_view const& input, thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), input.offsets_begin(), input.offsets_end(), - output_offset->mutable_view().begin(), + output_offset->mutable_view().begin(), [first = input.offsets_begin()] __device__(auto offset_index) { return offset_index - *first; }); diff --git a/cpp/src/lists/sequences.cu b/cpp/src/lists/sequences.cu index 55be5588bc7e..c57bd29bea25 100644 --- a/cpp/src/lists/sequences.cu +++ b/cpp/src/lists/sequences.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -156,7 +156,7 @@ std::unique_ptr sequences(column_view const& starts, // Generate list offsets for the output. auto [list_offsets, n_elements] = cudf::detail::make_offsets_child_column( sizes_input_it, sizes_input_it + sizes.size(), stream, mr); - auto const offsets_begin = list_offsets->view().template begin(); + auto const offsets_begin = list_offsets->view().template begin(); auto child = type_dispatcher(starts.type(), sequences_dispatcher{}, diff --git a/cpp/src/lists/stream_compaction/apply_boolean_mask.cu b/cpp/src/lists/stream_compaction/apply_boolean_mask.cu index 152fa5fc31da..04a9fa83a780 100644 --- a/cpp/src/lists/stream_compaction/apply_boolean_mask.cu +++ b/cpp/src/lists/stream_compaction/apply_boolean_mask.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -89,9 +89,9 @@ std::unique_ptr apply_mask(lists_column_view const& input, thrust::inclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), sizes_begin, sizes_end, - output_offsets_view.begin() + 1); - CUDF_CUDA_TRY(cudaMemsetAsync( - output_offsets_view.begin(), 0, sizeof(size_type), stream.value())); + output_offsets_view.begin() + 1); + CUDF_CUDA_TRY( + cudaMemsetAsync(output_offsets_view.begin(), 0, sizeof(int32_t), stream.value())); return output_offsets; }; diff --git a/cpp/src/lists/utilities.cu b/cpp/src/lists/utilities.cu index 365628e2cc8b..40b61da28181 100644 --- a/cpp/src/lists/utilities.cu +++ b/cpp/src/lists/utilities.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -32,10 +32,10 @@ std::unique_ptr reconstruct_offsets(column_view const& labels, { auto out_offsets = make_numeric_column( - data_type{type_to_id()}, n_lists + 1, mask_state::UNALLOCATED, stream, mr); + data_type{type_id::INT32}, n_lists + 1, mask_state::UNALLOCATED, stream, mr); auto const labels_begin = labels.template begin(); - auto const offsets_begin = out_offsets->mutable_view().template begin(); + auto const offsets_begin = out_offsets->mutable_view().template begin(); cudf::detail::labels_to_offsets(labels_begin, labels_begin + labels.size(), offsets_begin, @@ -50,15 +50,12 @@ std::unique_ptr get_normalized_offsets(lists_column_view const& input, { if (input.is_empty()) { return empty_like(input.offsets()); } - auto out_offsets = make_numeric_column(data_type(type_to_id()), - input.size() + 1, - cudf::mask_state::UNALLOCATED, - stream, - mr); + auto out_offsets = make_numeric_column( + data_type(type_id::INT32), input.size() + 1, cudf::mask_state::UNALLOCATED, stream, mr); thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), input.offsets_begin(), input.offsets_end(), - out_offsets->mutable_view().begin(), + out_offsets->mutable_view().begin(), [d_offsets = input.offsets_begin()] __device__(auto const offset_val) { // The first offset value, used for zero-normalizing offsets. return offset_val - *d_offsets; diff --git a/cpp/src/rolling/detail/rolling_collect_list.cu b/cpp/src/rolling/detail/rolling_collect_list.cu index 2b592cb60968..02e17abfabc6 100644 --- a/cpp/src/rolling/detail/rolling_collect_list.cu +++ b/cpp/src/rolling/detail/rolling_collect_list.cu @@ -42,8 +42,7 @@ std::unique_ptr get_list_child_to_list_row_mapping(cudf::column_view con // scatter result == [0, 0, 1, 0, 0, 2, 0, 0, 1, 0, 0, 1, 0] // - auto const num_child_rows{ - cudf::detail::get_value(offsets, offsets.size() - 1, stream)}; + auto const num_child_rows{cudf::detail::get_value(offsets, offsets.size() - 1, stream)}; auto per_row_mapping = make_fixed_width_column(data_type{type_to_id()}, num_child_rows, mask_state::UNALLOCATED, @@ -59,10 +58,10 @@ std::unique_ptr get_list_child_to_list_row_mapping(cudf::column_view con thrust::scatter_if(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), begin, begin + offsets.size() - 1, - offsets.begin(), + offsets.begin(), begin, // stencil iterator per_row_mapping_begin, - [offset = offsets.begin()] __device__(auto i) { + [offset = offsets.begin()] __device__(auto i) { return offset[i] != offset[i + 1]; }); // [0,0,1,0,0,3,...] @@ -141,7 +140,7 @@ std::pair, std::unique_ptr> purge_null_entries( new_sizes->mutable_view().template begin(), new_sizes->mutable_view().template end(), [d_gather_map = gather_map.template begin(), - d_old_offsets = offsets.template begin(), + d_old_offsets = offsets.template begin(), input_row_not_null] __device__(auto i) { return thrust::count_if(thrust::seq, d_gather_map + d_old_offsets[i], diff --git a/cpp/src/rolling/detail/rolling_collect_list.cuh b/cpp/src/rolling/detail/rolling_collect_list.cuh index f438ee68ce68..1ea94a2a76e6 100644 --- a/cpp/src/rolling/detail/rolling_collect_list.cuh +++ b/cpp/src/rolling/detail/rolling_collect_list.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -116,7 +116,7 @@ std::unique_ptr create_collect_gather_map(column_view const& child_offse gather_map->mutable_view().template begin(), cuda::proclaim_return_type( [d_offsets = - child_offsets.template begin(), // E.g. [0, 2, 5, 8, 11, 13] + child_offsets.template begin(), // E.g. [0, 2, 5, 8, 11, 13] d_groups = per_row_mapping.template begin(), // E.g. [0,0, 1,1,1, 2,2,2, 3,3,3, 4,4] d_prev = preceding_iter] __device__(auto i) { diff --git a/cpp/src/rolling/detail/rolling_operators.cuh b/cpp/src/rolling/detail/rolling_operators.cuh index f3d39f1edb06..f26150dc5986 100644 --- a/cpp/src/rolling/detail/rolling_operators.cuh +++ b/cpp/src/rolling/detail/rolling_operators.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -505,7 +505,7 @@ struct agg_specific_empty_output { if constexpr (op == aggregation::COLLECT_LIST) { return cudf::make_lists_column( - 0, make_empty_column(type_to_id()), empty_like(input), 0, {}); + 0, make_empty_column(type_id::INT32), empty_like(input), 0, {}); } return empty_like(input); diff --git a/cpp/src/strings/convert/convert_lists.cu b/cpp/src/strings/convert/convert_lists.cu index 8f27de09290c..ec3cfe23faa5 100644 --- a/cpp/src/strings/convert/convert_lists.cu +++ b/cpp/src/strings/convert/convert_lists.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -125,7 +125,7 @@ struct format_lists_fn { auto const view = get_nested_child(stack_idx); auto offsets = view.child(cudf::lists_column_view::offsets_column_index); - auto d_offsets = offsets.data() + view.offset(); + auto d_offsets = offsets.data() + view.offset(); // add pending separator if (item.separator == item_separator::LIST) { diff --git a/cpp/src/strings/repeat_strings.cu b/cpp/src/strings/repeat_strings.cu index 85ca5c7a2122..2faae9ea3b36 100644 --- a/cpp/src/strings/repeat_strings.cu +++ b/cpp/src/strings/repeat_strings.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -73,10 +73,10 @@ auto generate_empty_output(strings_column_view const& input, rmm::device_async_resource_ref mr) { auto offsets_column = make_numeric_column( - data_type{type_to_id()}, strings_count + 1, mask_state::UNALLOCATED, stream, mr); - CUDF_CUDA_TRY(cudaMemsetAsync(offsets_column->mutable_view().template data(), + data_type{type_id::INT32}, strings_count + 1, mask_state::UNALLOCATED, stream, mr); + CUDF_CUDA_TRY(cudaMemsetAsync(offsets_column->mutable_view().template data(), 0, - offsets_column->size() * sizeof(size_type), + offsets_column->size() * sizeof(int32_t), stream.value())); return make_strings_column(strings_count, diff --git a/cpp/src/transform/row_bit_count.cu b/cpp/src/transform/row_bit_count.cu index e604e47266f0..ba268ea58375 100644 --- a/cpp/src/transform/row_bit_count.cu +++ b/cpp/src/transform/row_bit_count.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -462,10 +462,9 @@ CUDF_KERNEL void compute_segment_sizes(device_span col // if this is a list column, update the working span from our offsets if (col.type().id() == type_id::LIST && col.size() > 0) { column_device_view const& offsets = col.child(lists_column_view::offsets_column_index); - auto const base_offset = offsets.data()[col.offset()]; - cur_span.row_start = - offsets.data()[cur_span.row_start + col.offset()] - base_offset; - cur_span.row_end = offsets.data()[cur_span.row_end + col.offset()] - base_offset; + auto const base_offset = offsets.data()[col.offset()]; + cur_span.row_start = offsets.data()[cur_span.row_start + col.offset()] - base_offset; + cur_span.row_end = offsets.data()[cur_span.row_end + col.offset()] - base_offset; } last_branch_depth = info[idx].branch_depth_end; diff --git a/cpp/tests/groupby/collect_list_tests.cpp b/cpp/tests/groupby/collect_list_tests.cpp index b416b22b6a86..c4e00b742796 100644 --- a/cpp/tests/groupby/collect_list_tests.cpp +++ b/cpp/tests/groupby/collect_list_tests.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -134,7 +134,7 @@ TYPED_TEST(groupby_collect_list_test, CollectOnEmptyInputLists) using LCW = cudf::test::lists_column_wrapper; - auto offsets = cudf::data_type{cudf::type_to_id()}; + auto offsets = cudf::data_type{cudf::type_id::INT32}; cudf::test::fixed_width_column_wrapper keys{}; auto values = @@ -162,30 +162,18 @@ TYPED_TEST(groupby_collect_list_test, CollectOnEmptyInputListsOfStructs) auto struct_child = LCW{}; auto struct_column = cudf::test::structs_column_wrapper{{struct_child}}; - auto values = - cudf::make_lists_column(0, - cudf::make_empty_column(cudf::type_to_id()), - struct_column.release(), - 0, - {}); + auto values = cudf::make_lists_column( + 0, cudf::make_empty_column(cudf::type_id::INT32), struct_column.release(), 0, {}); cudf::test::fixed_width_column_wrapper expect_keys{}; auto expect_struct_child = LCW{}; auto expect_struct_column = cudf::test::structs_column_wrapper{{expect_struct_child}}; - auto expect_child = - cudf::make_lists_column(0, - cudf::make_empty_column(cudf::type_to_id()), - expect_struct_column.release(), - 0, - {}); - auto expect_values = - cudf::make_lists_column(0, - cudf::make_empty_column(cudf::type_to_id()), - std::move(expect_child), - 0, - {}); + auto expect_child = cudf::make_lists_column( + 0, cudf::make_empty_column(cudf::type_id::INT32), expect_struct_column.release(), 0, {}); + auto expect_values = cudf::make_lists_column( + 0, cudf::make_empty_column(cudf::type_id::INT32), std::move(expect_child), 0, {}); auto agg = cudf::make_collect_list_aggregation(); test_single_agg(keys, values->view(), expect_keys, expect_values->view(), std::move(agg)); diff --git a/java/src/main/native/src/ColumnViewJni.cpp b/java/src/main/native/src/ColumnViewJni.cpp index b0d81a4a5ebc..3750f2ea880b 100644 --- a/java/src/main/native/src/ColumnViewJni.cpp +++ b/java/src/main/native/src/ColumnViewJni.cpp @@ -2298,7 +2298,7 @@ JNIEXPORT jlong JNICALL Java_ai_rapids_cudf_ColumnView_makeCudfColumnView(JNIEnv new cudf::column_view(cudf::data_type{cudf::type_id::STRING}, 0, nullptr, nullptr, 0)); } else { JNI_NULL_CHECK(env, j_offset, "offset is null", 0); - cudf::size_type* offsets = reinterpret_cast(j_offset); + int32_t* offsets = reinterpret_cast(j_offset); cudf::column_view offsets_column( cudf::data_type{cudf::type_id::INT32}, size + 1, offsets, nullptr, 0); return ptr_as_jlong(new cudf::column_view(cudf::data_type{cudf::type_id::STRING}, @@ -2314,11 +2314,11 @@ JNIEXPORT jlong JNICALL Java_ai_rapids_cudf_ColumnView_makeCudfColumnView(JNIEnv cudf::jni::native_jpointerArray children(env, j_children); JNI_ARG_CHECK(env, (children.size() == 1), "LIST children size is not 1", 0); cudf::size_type offsets_size = 0; - cudf::size_type* offsets = nullptr; + int32_t* offsets = nullptr; if (size != 0) { JNI_NULL_CHECK(env, j_offset, "offset is null", 0); offsets_size = size + 1; - offsets = reinterpret_cast(j_offset); + offsets = reinterpret_cast(j_offset); } cudf::column_view offsets_column( cudf::data_type{cudf::type_id::INT32}, offsets_size, offsets, nullptr, 0); diff --git a/java/src/main/native/src/ColumnViewJni.cu b/java/src/main/native/src/ColumnViewJni.cu index a140dac9501e..ff58835591ae 100644 --- a/java/src/main/native/src/ColumnViewJni.cu +++ b/java/src/main/native/src/ColumnViewJni.cu @@ -179,8 +179,8 @@ std::unique_ptr lists_distinct_by_key(cudf::lists_column_view cons // Assemble a lists column of structs. auto out_offsets = make_numeric_column( - data_type{type_to_id()}, input.size() + 1, mask_state::UNALLOCATED, stream); - auto const offsets_begin = out_offsets->mutable_view().template begin(); + data_type{type_id::INT32}, input.size() + 1, mask_state::UNALLOCATED, stream); + auto const offsets_begin = out_offsets->mutable_view().template begin(); auto const labels_begin = out_labels.template begin(); cudf::detail::labels_to_offsets(labels_begin, labels_begin + out_labels.size(), From b7f310a6f08be3fd18ffc6846880f136eb437b99 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 10 Aug 2026 14:10:19 -0400 Subject: [PATCH 2/3] fix additional size-type=int32 assumptions --- cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh | 9 +++++++-- cpp/include/cudf/lists/detail/scatter.cuh | 4 ++-- cpp/src/io/json/host_tree_algorithms.cu | 2 +- cpp/src/io/json/parser_features.cpp | 2 +- cpp/src/lists/combine/concatenate_list_elements.cu | 2 +- cpp/src/lists/copying/concatenate.cu | 2 +- cpp/src/lists/copying/copying.cu | 2 +- cpp/src/lists/sequences.cu | 6 +++--- java/src/main/native/src/ColumnViewJni.cu | 4 ++-- 9 files changed, 19 insertions(+), 14 deletions(-) diff --git a/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh b/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh index 14fff8857154..b19cc1c7a9af 100644 --- a/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh +++ b/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh @@ -317,9 +317,14 @@ std::pair, size_type> make_offsets_child_column( // using exclusive-scan technically requires count+1 input values even though // the final input value is never used. // The input iterator is wrapped here to allow the last value to be safely read. + // The input sizes are deliberately not narrowed to the 32-bit offsets type here -- doing so + // would corrupt individual sizes larger than `int32_t` before they reach the scan, so the + // overflow check below could no longer detect the overflow. Narrowing happens only on write + // to `d_offsets`, after the accumulated total has been validated. + using SizeType = cuda::std::iter_value_t; auto map_fn = - cuda::proclaim_return_type([begin, count] __device__(size_type idx) -> int32_t { - return idx < count ? static_cast(begin[idx]) : int32_t{0}; + cuda::proclaim_return_type([begin, count] __device__(size_type idx) -> SizeType { + return idx < count ? begin[idx] : SizeType{0}; }); auto input_itr = cudf::detail::make_counting_transform_iterator(0, map_fn); // Use the sizes-to-offsets iterator to compute the total number of elements diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index c84c8f67bcfc..d141adfaf973 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -235,8 +235,8 @@ std::unique_ptr scatter(scalar const& slr, thrust::sequence(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), offset_column->mutable_view().begin(), offset_column->mutable_view().end(), - 0, - lv->view().size()); + int32_t{0}, + static_cast(lv->view().size())); auto wrapped = column_view(data_type{type_id::LIST}, 1, nullptr, diff --git a/cpp/src/io/json/host_tree_algorithms.cu b/cpp/src/io/json/host_tree_algorithms.cu index a85b996ac55f..6526b7089e7d 100644 --- a/cpp/src/io/json/host_tree_algorithms.cu +++ b/cpp/src/io/json/host_tree_algorithms.cu @@ -1154,7 +1154,7 @@ void scatter_offsets(tree_meta_t const& tree, col.child_offsets.begin(), col.child_offsets.end(), col.child_offsets.begin(), - cuda::maximum{}); + cuda::maximum{}); } } stream.synchronize(); diff --git a/cpp/src/io/json/parser_features.cpp b/cpp/src/io/json/parser_features.cpp index 1ee13c1f747d..a502fed7897a 100644 --- a/cpp/src/io/json/parser_features.cpp +++ b/cpp/src/io/json/parser_features.cpp @@ -117,7 +117,7 @@ struct allnull_column_functor { [[nodiscard]] auto make_zeroed_offsets(size_type size) const { auto offsets_buff = - cudf::detail::make_zeroed_device_uvector_async(size + 1, stream, mr); + cudf::detail::make_zeroed_device_uvector_async(size + 1, stream, mr); return std::make_unique(std::move(offsets_buff), rmm::device_buffer{}, 0); } diff --git a/cpp/src/lists/combine/concatenate_list_elements.cu b/cpp/src/lists/combine/concatenate_list_elements.cu index ebcc183437e4..376845d31b31 100644 --- a/cpp/src/lists/combine/concatenate_list_elements.cu +++ b/cpp/src/lists/combine/concatenate_list_elements.cu @@ -257,7 +257,7 @@ std::unique_ptr concatenate_list_elements(column_view const& input, if (child.size() == 0) { auto const num_rows = input.size(); auto out_offsets = cudf::make_column_from_scalar( - numeric_scalar(0, true, stream, mr), num_rows + 1, stream, mr); + numeric_scalar(0, true, stream, mr), num_rows + 1, stream, mr); // Use the grandchild's schema (not child's) so out_entries has type T, not list. auto out_entries = cudf::empty_like(lists_column_view(child).child()); return make_lists_column(num_rows, diff --git a/cpp/src/lists/copying/concatenate.cu b/cpp/src/lists/copying/concatenate.cu index 3324f18ec3c2..4ed2b30d1e5a 100644 --- a/cpp/src/lists/copying/concatenate.cu +++ b/cpp/src/lists/copying/concatenate.cu @@ -70,7 +70,7 @@ std::unique_ptr merge_offsets(host_span columns offsets.begin() + c.offset(), offsets.begin() + c.offset() + c.size() + 1, d_merged_offsets.begin() + count, - [local_shift] __device__(size_type offset) { return offset + local_shift; }); + [local_shift] __device__(int32_t offset) -> int32_t { return offset + local_shift; }); shift += c.get_sliced_child(stream).size(); count += c.size(); diff --git a/cpp/src/lists/copying/copying.cu b/cpp/src/lists/copying/copying.cu index a4bcf7688bdd..e25a99230047 100644 --- a/cpp/src/lists/copying/copying.cu +++ b/cpp/src/lists/copying/copying.cu @@ -44,7 +44,7 @@ std::unique_ptr copy_slice(lists_column_view const& lists, auto start_offset = cudf::detail::get_value(lists.offsets(), start, stream); auto end_offset = cudf::detail::get_value(lists.offsets(), end, stream); - rmm::device_uvector out_offsets(offsets_count, stream); + rmm::device_uvector out_offsets(offsets_count, stream); // Compute the offsets column of the result: thrust::transform( diff --git a/cpp/src/lists/sequences.cu b/cpp/src/lists/sequences.cu index c57bd29bea25..2d75f0e0f146 100644 --- a/cpp/src/lists/sequences.cu +++ b/cpp/src/lists/sequences.cu @@ -39,7 +39,7 @@ struct tabulator { T const* const starts; T const* const steps; - size_type const* const offsets; + int32_t const* const offsets; template static T __device__ multiply(U x, size_type times) @@ -80,7 +80,7 @@ struct sequences_dispatcher { size_type n_elements, column_view const& starts, std::optional const& steps, - size_type const* offsets, + int32_t const* offsets, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { @@ -100,7 +100,7 @@ struct sequences_functor()>> { size_type n_elements, column_view const& starts, std::optional const& steps, - size_type const* offsets, + int32_t const* offsets, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { diff --git a/java/src/main/native/src/ColumnViewJni.cu b/java/src/main/native/src/ColumnViewJni.cu index ff58835591ae..028e3f5a548d 100644 --- a/java/src/main/native/src/ColumnViewJni.cu +++ b/java/src/main/native/src/ColumnViewJni.cu @@ -38,8 +38,8 @@ std::unique_ptr generate_list_offsets(cudf::column_view const& lis CUDF_EXPECTS(list_length.type().id() == cudf::type_id::INT32, "Input column does not have type INT32."); - auto const begin_iter = list_length.template begin(); - auto const end_iter = list_length.template end(); + auto const begin_iter = list_length.template begin(); + auto const end_iter = list_length.template end(); auto offsets_column = make_numeric_column( data_type{type_id::INT32}, list_length.size() + 1, mask_state::UNALLOCATED, stream); From 3d4ca87c12ccf7529527f0c0687e7427725f2a09 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 10 Aug 2026 14:46:26 -0400 Subject: [PATCH 3/3] tdigest has list columns --- cpp/src/quantiles/tdigest/tdigest.cu | 16 ++++----- .../quantiles/tdigest/tdigest_aggregation.cu | 36 +++++++++---------- cpp/src/quantiles/tdigest/tdigest_util.cuh | 4 +-- cpp/src/strings/search/find_multiple.cu | 6 ++-- cpp/src/text/minhash.cu | 2 +- cpp/tests/groupby/collect_list_tests.cpp | 28 ++++++++++----- 6 files changed, 52 insertions(+), 40 deletions(-) diff --git a/cpp/src/quantiles/tdigest/tdigest.cu b/cpp/src/quantiles/tdigest/tdigest.cu index bf793aecb3bc..cf6eaaf8654f 100644 --- a/cpp/src/quantiles/tdigest/tdigest.cu +++ b/cpp/src/quantiles/tdigest/tdigest.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -59,7 +59,7 @@ struct make_centroid { // kernel for computing percentiles on input tdigest (mean, weight) centroid data. template -CUDF_KERNEL void compute_percentiles_kernel(device_span tdigest_offsets, +CUDF_KERNEL void compute_percentiles_kernel(device_span tdigest_offsets, column_device_view percentiles, CentroidIter centroids_, double const* min_, @@ -193,8 +193,8 @@ std::unique_ptr compute_approx_percentiles(tdigest_column_view const& in auto keys = cudf::detail::make_counting_transform_iterator( 0, cuda::proclaim_return_type( - [offsets_begin = offsets.begin(), - offsets_end = offsets.end()] __device__(size_type i) { + [offsets_begin = offsets.begin(), + offsets_end = offsets.end()] __device__(size_type i) { return cuda::std::distance( offsets_begin, cuda::std::prev(thrust::upper_bound(thrust::seq, offsets_begin, offsets_end, i))); @@ -234,7 +234,7 @@ std::unique_ptr compute_approx_percentiles(tdigest_column_view const& in constexpr size_type block_size = 256; cudf::detail::grid_1d const grid(percentiles.size() * input.size(), block_size); compute_percentiles_kernel<<>>( - {offsets.begin(), static_cast(offsets.size())}, + {offsets.begin(), static_cast(offsets.size())}, *percentiles_cdv, centroids, tdv.min_begin(), @@ -291,8 +291,8 @@ std::unique_ptr make_empty_tdigests_column(size_type num_rows, auto offsets = cudf::make_fixed_width_column( data_type(type_id::INT32), num_rows + 1, mask_state::UNALLOCATED, stream, mr); thrust::fill(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - offsets->mutable_view().begin(), - offsets->mutable_view().end(), + offsets->mutable_view().begin(), + offsets->mutable_view().end(), 0); auto min_col = cudf::make_numeric_column( @@ -359,7 +359,7 @@ std::unique_ptr percentile_approx(tdigest_column_view const& input, thrust::exclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), row_size_iter, row_size_iter + input.size() + 1, - offsets->mutable_view().begin()); + offsets->mutable_view().begin()); if (percentiles.size() == 0 || all_empty_rows) { return cudf::make_lists_column( diff --git a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu index ffa3a3d7a28a..145b4a1d2160 100644 --- a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu +++ b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu @@ -184,7 +184,7 @@ template struct nearest_value_centroid_weights { double const* cumulative_weights; // cumulative weights of non-empty clusters GroupOffsetsIter group_offsets; // groups - size_type const* tdigest_offsets; // tdigests within a group + int32_t const* tdigest_offsets; // tdigests within a group cuda::std::pair operator() CUDF_HOST_DEVICE(double next_limit, size_type group_index) const @@ -258,7 +258,7 @@ struct cumulative_centroid_weight { GroupLabelsIter group_labels; // group labels for each tdigest including empty ones GroupOffsetsIter group_offsets; // groups // Host-device span, as the offsets may reside in either device memory or pinned host memory - cuda::std::span tdigest_offsets; // tdigests with a group + cuda::std::span tdigest_offsets; // tdigests with a group /** * @brief Returns the cumulative weight for a given value index. The index `n` is the index of @@ -313,7 +313,7 @@ template struct centroid_group_info { double const* cumulative_weights; // cumulative weights of non-empty clusters GroupOffsetsIter group_offsets; - size_type const* tdigest_offsets; + int32_t const* tdigest_offsets; CUDF_HOST_DEVICE cuda::std::tuple operator()( size_type group_index) const @@ -1297,7 +1297,7 @@ struct typed_reduce_tdigest { template struct group_num_clusters_func { GroupOffsetsIter group_offsets; - size_type const* tdigest_offsets; + int32_t const* tdigest_offsets; __device__ size_type operator()(size_type group_index) { @@ -1324,7 +1324,7 @@ struct group_is_empty { template struct group_key_func { GroupLabelsIter group_labels; - size_type const* tdigest_offsets; + int32_t const* tdigest_offsets; size_type num_tdigest_offsets; /** @@ -1382,7 +1382,7 @@ std::pair, rmm::device_uvector> generate_mer auto centroid_offsets = cudf::detail::make_counting_transform_iterator( 0, cuda::proclaim_return_type( - [group_offsets, tdigest_offsets = tdv.centroids().offsets().begin()] __device__( + [group_offsets, tdigest_offsets = tdv.centroids().offsets().begin()] __device__( size_type i) { return tdigest_offsets[group_offsets[i]]; })); // perform the sort using the means as the key @@ -1484,7 +1484,7 @@ std::unique_ptr merge_tdigests(tdigest_column_view const& tdv, auto group_num_clusters = cudf::detail::make_counting_transform_iterator( 0, group_num_clusters_func{group_offsets, - tdigest_offsets.begin()}); + tdigest_offsets.begin()}); thrust::replace_if(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), merged_min_col->mutable_view().begin(), merged_min_col->mutable_view().end(), @@ -1517,7 +1517,7 @@ std::unique_ptr merge_tdigests(tdigest_column_view const& tdv, iter + num_centroids, group_keys.begin(), group_key_func{ - group_labels, tdigest_offsets.begin(), tdigest_offsets.size()}); + group_labels, tdigest_offsets.begin(), tdigest_offsets.size()}); thrust::inclusive_scan_by_key( rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), group_keys.begin(), @@ -1547,10 +1547,10 @@ std::unique_ptr merge_tdigests(tdigest_column_view const& tdv, rmm::device_uvector p_cumulative_weights(cumulative_weights, stream, pinned_mr); - rmm::device_uvector p_tdigest_offsets(tdigest_offsets.size(), stream, pinned_mr); + rmm::device_uvector p_tdigest_offsets(tdigest_offsets.size(), stream, pinned_mr); thrust::copy(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - tdigest_offsets.begin(), - tdigest_offsets.begin() + p_tdigest_offsets.size(), + tdigest_offsets.begin(), + tdigest_offsets.begin() + p_tdigest_offsets.size(), p_tdigest_offsets.begin()); rmm::device_uvector _p_group_labels(num_group_labels, stream, pinned_mr); @@ -1572,7 +1572,7 @@ std::unique_ptr merge_tdigests(tdigest_column_view const& tdv, p_cumulative_weights.begin(), p_group_labels, p_group_offsets, - cuda::std::span{p_tdigest_offsets.begin(), p_tdigest_offsets.size()}}, + cuda::std::span{p_tdigest_offsets.begin(), p_tdigest_offsets.size()}}, has_nulls, stream, mr); @@ -1583,15 +1583,15 @@ std::unique_ptr merge_tdigests(tdigest_column_view const& tdv, delta, num_groups, nearest_value_centroid_weights{ - cumulative_weights.begin(), group_offsets, tdigest_offsets.begin()}, + cumulative_weights.begin(), group_offsets, tdigest_offsets.begin()}, centroid_group_info{ - cumulative_weights.begin(), group_offsets, tdigest_offsets.begin()}, + cumulative_weights.begin(), group_offsets, tdigest_offsets.begin()}, cumulative_centroid_weight{ cumulative_weights.begin(), group_labels, group_offsets, - cuda::std::span{tdigest_offsets.begin(), - static_cast(tdigest_offsets.size())}}, + cuda::std::span{tdigest_offsets.begin(), + static_cast(tdigest_offsets.size())}}, has_nulls, stream, mr); @@ -1610,8 +1610,8 @@ std::unique_ptr merge_tdigests(tdigest_column_view const& tdv, cumulative_weights.begin(), group_labels, group_offsets, - cuda::std::span{tdigest_offsets.begin(), - static_cast(tdigest_offsets.size())}}, + cuda::std::span{tdigest_offsets.begin(), + static_cast(tdigest_offsets.size())}}, std::move(merged_min_col), std::move(merged_max_col), cinfo, diff --git a/cpp/src/quantiles/tdigest/tdigest_util.cuh b/cpp/src/quantiles/tdigest/tdigest_util.cuh index 7e7557248f41..fc6034a67a56 100644 --- a/cpp/src/quantiles/tdigest/tdigest_util.cuh +++ b/cpp/src/quantiles/tdigest/tdigest_util.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2022, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -15,7 +15,7 @@ namespace detail { * @brief Functor to compute the size of each tdigest of a column */ struct tdigest_size_fn { - size_type const* offsets; ///< Offsets of the t-digest column + int32_t const* offsets; ///< Offsets of the t-digest column /** * @brief Returns size of the each tdigest in the column * diff --git a/cpp/src/strings/search/find_multiple.cu b/cpp/src/strings/search/find_multiple.cu index 198528f11ef6..3d081c3e6e77 100644 --- a/cpp/src/strings/search/find_multiple.cu +++ b/cpp/src/strings/search/find_multiple.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -69,8 +69,8 @@ std::unique_ptr find_multiple(strings_column_view const& input, auto offsets = cudf::detail::sequence( strings_count + 1, - numeric_scalar(0, true, stream, cudf::get_current_device_resource_ref()), - numeric_scalar(targets_count, true, stream, cudf::get_current_device_resource_ref()), + numeric_scalar(0, true, stream, cudf::get_current_device_resource_ref()), + numeric_scalar(targets_count, true, stream, cudf::get_current_device_resource_ref()), stream, mr); return make_lists_column( diff --git a/cpp/src/text/minhash.cu b/cpp/src/text/minhash.cu index 9669a196057b..80154e15ce3c 100644 --- a/cpp/src/text/minhash.cu +++ b/cpp/src/text/minhash.cu @@ -194,7 +194,7 @@ CUDF_KERNEL void minhash_ngrams_kernel(cudf::lists_column_device_view const d_in if (d_input.is_null(row_idx)) { return; } // retrieve this row's offset to locate the output position in d_hashes - auto const offsets_itr = d_input.offsets().data() + d_input.offset(); + auto const offsets_itr = d_input.offsets().data() + d_input.offset(); auto const offset = offsets_itr[row_idx]; auto const size_row = offsets_itr[row_idx + 1] - offset; if (size_row == 0) { return; } diff --git a/cpp/tests/groupby/collect_list_tests.cpp b/cpp/tests/groupby/collect_list_tests.cpp index c4e00b742796..b416b22b6a86 100644 --- a/cpp/tests/groupby/collect_list_tests.cpp +++ b/cpp/tests/groupby/collect_list_tests.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2021-2024, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -134,7 +134,7 @@ TYPED_TEST(groupby_collect_list_test, CollectOnEmptyInputLists) using LCW = cudf::test::lists_column_wrapper; - auto offsets = cudf::data_type{cudf::type_id::INT32}; + auto offsets = cudf::data_type{cudf::type_to_id()}; cudf::test::fixed_width_column_wrapper keys{}; auto values = @@ -162,18 +162,30 @@ TYPED_TEST(groupby_collect_list_test, CollectOnEmptyInputListsOfStructs) auto struct_child = LCW{}; auto struct_column = cudf::test::structs_column_wrapper{{struct_child}}; - auto values = cudf::make_lists_column( - 0, cudf::make_empty_column(cudf::type_id::INT32), struct_column.release(), 0, {}); + auto values = + cudf::make_lists_column(0, + cudf::make_empty_column(cudf::type_to_id()), + struct_column.release(), + 0, + {}); cudf::test::fixed_width_column_wrapper expect_keys{}; auto expect_struct_child = LCW{}; auto expect_struct_column = cudf::test::structs_column_wrapper{{expect_struct_child}}; - auto expect_child = cudf::make_lists_column( - 0, cudf::make_empty_column(cudf::type_id::INT32), expect_struct_column.release(), 0, {}); - auto expect_values = cudf::make_lists_column( - 0, cudf::make_empty_column(cudf::type_id::INT32), std::move(expect_child), 0, {}); + auto expect_child = + cudf::make_lists_column(0, + cudf::make_empty_column(cudf::type_to_id()), + expect_struct_column.release(), + 0, + {}); + auto expect_values = + cudf::make_lists_column(0, + cudf::make_empty_column(cudf::type_to_id()), + std::move(expect_child), + 0, + {}); auto agg = cudf::make_collect_list_aggregation(); test_single_agg(keys, values->view(), expect_keys, expect_values->view(), std::move(agg));