From 4f002fc381e9bd6eab88bd20c3db6fb96aaeb996 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 20 Feb 2024 14:20:22 +0000 Subject: [PATCH 1/2] Replace `thrust::tuple` with `cuda::std::tuple` We want to move away from thrust types in API boundaries in favor of the more suited `cuda::std` types --- cpp/benchmarks/common/generate_input.cu | 14 ++--- cpp/benchmarks/string/url_decode.cu | 8 +-- cpp/include/cudf/detail/merge.hpp | 4 +- cpp/include/cudf/detail/null_mask.cuh | 14 ++--- cpp/include/cudf/detail/replace/nulls.cuh | 6 +- cpp/include/cudf/strings/detail/merge.cuh | 4 +- .../detail/strings_column_factories.cuh | 8 +-- cpp/include/cudf_test/tdigest_utilities.cuh | 4 +- cpp/src/copying/contiguous_split.cu | 6 +- cpp/src/groupby/sort/group_correlation.cu | 6 +- cpp/src/groupby/sort/group_merge_m2.cu | 10 ++-- cpp/src/groupby/sort/group_replace_nulls.cu | 6 +- cpp/src/io/comp/nvcomp_adapter.cu | 6 +- cpp/src/io/comp/statistics.cu | 4 +- cpp/src/io/json/json_column.cu | 10 ++-- cpp/src/io/json/json_tree.cu | 4 +- cpp/src/io/json/nested_json_gpu.cu | 20 +++---- cpp/src/io/json/write_json.cu | 4 +- cpp/src/io/orc/stripe_enc.cu | 2 +- cpp/src/io/parquet/page_enc.cu | 2 +- cpp/src/io/parquet/page_hdr.cu | 26 ++++---- cpp/src/io/parquet/reader_impl_helpers.hpp | 2 +- cpp/src/io/text/bgzip_data_chunk_source.cu | 14 ++--- cpp/src/io/utilities/data_casting.cu | 8 +-- cpp/src/io/utilities/string_parsing.hpp | 6 +- cpp/src/io/utilities/type_inference.cu | 10 ++-- cpp/src/join/hash_join.cu | 10 ++-- cpp/src/json/json_path.cu | 2 +- cpp/src/lists/dremel.cu | 12 ++-- cpp/src/merge/merge.cu | 6 +- cpp/src/partitioning/round_robin.cu | 2 +- .../quantiles/tdigest/tdigest_aggregation.cu | 60 +++++++++---------- cpp/src/reductions/histogram.cu | 8 +-- cpp/src/replace/clamp.cu | 12 ++-- cpp/src/replace/nulls.cu | 6 +- cpp/src/replace/replace.cu | 2 +- cpp/src/sort/rank.cu | 6 +- cpp/src/strings/split/split_re.cu | 2 +- cpp/src/text/minhash.cu | 2 +- cpp/tests/io/json_type_cast_test.cu | 14 ++--- cpp/tests/io/nested_json_test.cpp | 6 +- cpp/tests/io/type_inference_test.cu | 16 ++--- cpp/tests/reductions/scan_tests.cpp | 12 ++-- cpp/tests/utilities/tdigest_utilities.cu | 8 +-- 44 files changed, 197 insertions(+), 197 deletions(-) diff --git a/cpp/benchmarks/common/generate_input.cu b/cpp/benchmarks/common/generate_input.cu index 0ea13957868..567cabf4aa0 100644 --- a/cpp/benchmarks/common/generate_input.cu +++ b/cpp/benchmarks/common/generate_input.cu @@ -51,7 +51,7 @@ #include #include #include -#include +#include #include @@ -464,9 +464,9 @@ std::unique_ptr create_random_column(data_profile const& profile, struct valid_or_zero { template - __device__ T operator()(thrust::tuple len_valid) const + __device__ T operator()(cuda::std::tuple len_valid) const { - return thrust::get<1>(len_valid) ? thrust::get<0>(len_valid) : T{0}; + return cuda::std::get<1>(len_valid) ? cuda::std::get<0>(len_valid) : T{0}; } }; @@ -481,10 +481,10 @@ struct string_generator { // range 32-127 is ASCII; 127-136 will be multi-byte UTF-8 { } - __device__ void operator()(thrust::tuple str_begin_end) + __device__ void operator()(cuda::std::tuple str_begin_end) { - auto begin = thrust::get<0>(str_begin_end); - auto end = thrust::get<1>(str_begin_end); + auto begin = cuda::std::get<0>(str_begin_end); + auto end = cuda::std::get<1>(str_begin_end); engine.discard(begin); for (auto i = begin; i < end; ++i) { auto ch = char_dist(engine); @@ -519,7 +519,7 @@ std::unique_ptr create_random_utf8_string_column(data_profile cons cuda::proclaim_return_type([] __device__(auto) { return 0; }), thrust::logical_not{}); auto valid_lengths = thrust::make_transform_iterator( - thrust::make_zip_iterator(thrust::make_tuple(lengths.begin(), null_mask.begin())), + thrust::make_zip_iterator(cuda::std::make_tuple(lengths.begin(), null_mask.begin())), valid_or_zero{}); rmm::device_uvector offsets(num_rows + 1, cudf::get_default_stream()); thrust::exclusive_scan( diff --git a/cpp/benchmarks/string/url_decode.cu b/cpp/benchmarks/string/url_decode.cu index b3aeb69e5ea..36d1c7e8ace 100644 --- a/cpp/benchmarks/string/url_decode.cu +++ b/cpp/benchmarks/string/url_decode.cu @@ -31,7 +31,7 @@ #include #include #include -#include +#include struct url_string_generator { char* chars; @@ -43,10 +43,10 @@ struct url_string_generator { { } - __device__ void operator()(thrust::tuple str_begin_end) + __device__ void operator()(cuda::std::tuple str_begin_end) { - auto begin = thrust::get<0>(str_begin_end); - auto end = thrust::get<1>(str_begin_end); + auto begin = cuda::std::get<0>(str_begin_end); + auto end = cuda::std::get<1>(str_begin_end); engine.discard(begin); for (auto i = begin; i < end; ++i) { if (esc_seq_dist(engine) < esc_seq_chance and i < end - 3) { diff --git a/cpp/include/cudf/detail/merge.hpp b/cpp/include/cudf/detail/merge.hpp index 2167a484214..27f94d085d1 100644 --- a/cpp/include/cudf/detail/merge.hpp +++ b/cpp/include/cudf/detail/merge.hpp @@ -29,8 +29,8 @@ namespace detail { enum class side : bool { LEFT, RIGHT }; /** - * @brief Tagged index type: `thrust::get<0>` indicates left/right side, - * `thrust::get<1>` indicates the row index + * @brief Tagged index type: `cuda::std::get<0>` indicates left/right side, + * `cuda::std::get<1>` indicates the row index */ using index_type = thrust::pair; diff --git a/cpp/include/cudf/detail/null_mask.cuh b/cpp/include/cudf/detail/null_mask.cuh index e57d85f2998..f62679abac5 100644 --- a/cpp/include/cudf/detail/null_mask.cuh +++ b/cpp/include/cudf/detail/null_mask.cuh @@ -35,7 +35,7 @@ #include #include #include -#include +#include #include @@ -334,8 +334,8 @@ rmm::device_uvector segmented_count_bits(bitmask_type const* bitmask, thrust::make_zip_iterator(first_bit_indices_begin, last_bit_indices_begin); auto segment_length_iterator = thrust::transform_iterator( segments_begin, cuda::proclaim_return_type([] __device__(auto const& segment) { - auto const begin = thrust::get<0>(segment); - auto const end = thrust::get<1>(segment); + auto const begin = cuda::std::get<0>(segment); + auto const end = cuda::std::get<1>(segment); return end - begin; })); thrust::transform(rmm::exec_policy(stream), @@ -546,8 +546,8 @@ std::pair segmented_null_mask_reduction( thrust::make_zip_iterator(first_bit_indices_begin, last_bit_indices_begin); auto const segment_length_iterator = thrust::make_transform_iterator( segments_begin, cuda::proclaim_return_type([] __device__(auto const& segment) { - auto const begin = thrust::get<0>(segment); - auto const end = thrust::get<1>(segment); + auto const begin = cuda::std::get<0>(segment); + auto const end = cuda::std::get<1>(segment); return end - begin; })); @@ -579,8 +579,8 @@ std::pair segmented_null_mask_reduction( length_and_valid_count, length_and_valid_count + num_segments, [null_handling, valid_initial_value] __device__(auto const& length_and_valid_count) { - auto const length = thrust::get<0>(length_and_valid_count); - auto const valid_count = thrust::get<1>(length_and_valid_count); + auto const length = cuda::std::get<0>(length_and_valid_count); + auto const valid_count = cuda::std::get<1>(length_and_valid_count); return (null_handling == null_policy::EXCLUDE) ? (valid_initial_value.value_or(false) || valid_count > 0) : (valid_initial_value.value_or(length > 0) && valid_count == length); diff --git a/cpp/include/cudf/detail/replace/nulls.cuh b/cpp/include/cudf/detail/replace/nulls.cuh index d691ef5ce8e..f5061ee5a9a 100644 --- a/cpp/include/cudf/detail/replace/nulls.cuh +++ b/cpp/include/cudf/detail/replace/nulls.cuh @@ -18,12 +18,12 @@ #include #include -#include +#include namespace cudf { namespace detail { -using idx_valid_pair_t = thrust::tuple; +using idx_valid_pair_t = cuda::std::tuple; /** * @brief Functor used by `replace_nulls(replace_policy)` to determine the index to gather from in @@ -36,7 +36,7 @@ using idx_valid_pair_t = thrust::tuple; struct replace_policy_functor { __device__ idx_valid_pair_t operator()(idx_valid_pair_t const& lhs, idx_valid_pair_t const& rhs) { - return thrust::get<1>(rhs) ? rhs : lhs; + return cuda::std::get<1>(rhs) ? rhs : lhs; } }; diff --git a/cpp/include/cudf/strings/detail/merge.cuh b/cpp/include/cudf/strings/detail/merge.cuh index 8049895c3c2..bc40c5472e7 100644 --- a/cpp/include/cudf/strings/detail/merge.cuh +++ b/cpp/include/cudf/strings/detail/merge.cuh @@ -30,7 +30,7 @@ #include #include #include -#include +#include #include @@ -42,7 +42,7 @@ namespace detail { * * Caller must set the validity mask in the output column. * - * @tparam row_order_iterator This must be an iterator for type thrust::tuple. + * @tparam row_order_iterator This must be an iterator for type cuda::std::tuple. * * @param lhs First column. * @param rhs Second column. diff --git a/cpp/include/cudf/strings/detail/strings_column_factories.cuh b/cpp/include/cudf/strings/detail/strings_column_factories.cuh index 0adf6e362be..8884a5fca32 100644 --- a/cpp/include/cudf/strings/detail/strings_column_factories.cuh +++ b/cpp/include/cudf/strings/detail/strings_column_factories.cuh @@ -35,7 +35,7 @@ #include #include #include -#include +#include #include @@ -120,13 +120,13 @@ std::unique_ptr make_strings_column(IndexPairIterator begin, auto chars_data = rmm::device_uvector(bytes, stream, mr); auto d_chars = chars_data.data(); auto copy_chars = [d_chars] __device__(auto item) { - string_index_pair const str = thrust::get<0>(item); - size_type const offset = thrust::get<1>(item); + string_index_pair const str = cuda::std::get<0>(item); + size_type const offset = cuda::std::get<1>(item); if (str.first != nullptr) memcpy(d_chars + offset, str.first, str.second); }; thrust::for_each_n(rmm::exec_policy(stream), thrust::make_zip_iterator( - thrust::make_tuple(begin, offsets_view.template begin())), + cuda::std::make_tuple(begin, offsets_view.template begin())), strings_count, copy_chars); return chars_data; diff --git a/cpp/include/cudf_test/tdigest_utilities.cuh b/cpp/include/cudf_test/tdigest_utilities.cuh index 742cd764a1f..1b68262c38d 100644 --- a/cpp/include/cudf_test/tdigest_utilities.cuh +++ b/cpp/include/cudf_test/tdigest_utilities.cuh @@ -33,14 +33,14 @@ #include #include #include -#include +#include // for use with groupby and reduction aggregation tests. namespace cudf { namespace test { -using expected_value = thrust::tuple; +using expected_value = cuda::std::tuple; /** * @brief Device functor to compute min of a sequence of values serially. diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index d711f40605a..3bf105440c8 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -43,7 +43,7 @@ #include #include #include -#include +#include #include @@ -947,7 +947,7 @@ struct size_of_helper { */ struct num_batches_func { thrust::pair const* const batches; - __device__ std::size_t operator()(size_type i) const { return thrust::get<0>(batches[i]); } + __device__ std::size_t operator()(size_type i) const { return cuda::std::get<0>(batches[i]); } }; /** @@ -1466,7 +1466,7 @@ std::unique_ptr chunk_iteration_state::create( out_to_in_index] __device__(size_type i) { size_type const in_buf_index = out_to_in_index(i); size_type const batch_index = i - d_batch_offsets[in_buf_index]; - auto const batch_size = thrust::get<1>(batches[in_buf_index]); + auto const batch_size = cuda::std::get<1>(batches[in_buf_index]); dst_buf_info const& in = d_orig_dst_buf_info[in_buf_index]; // adjust info diff --git a/cpp/src/groupby/sort/group_correlation.cu b/cpp/src/groupby/sort/group_correlation.cu index 887e82e66df..0b1691dcc27 100644 --- a/cpp/src/groupby/sort/group_correlation.cu +++ b/cpp/src/groupby/sort/group_correlation.cu @@ -33,7 +33,7 @@ #include #include #include -#include +#include #include @@ -187,7 +187,7 @@ std::unique_ptr group_correlation(column_view const& covariance, CUDF_EXPECTS(covariance.type().id() == type_id::FLOAT64, "Covariance result must be FLOAT64"); auto stddev0_ptr = stddev_0.begin(); auto stddev1_ptr = stddev_1.begin(); - auto stddev_iter = thrust::make_zip_iterator(thrust::make_tuple(stddev0_ptr, stddev1_ptr)); + auto stddev_iter = thrust::make_zip_iterator(cuda::std::make_tuple(stddev0_ptr, stddev1_ptr)); auto result = make_numeric_column(covariance.type(), covariance.size(), cudf::detail::copy_bitmask(covariance, stream, mr), @@ -201,7 +201,7 @@ std::unique_ptr group_correlation(column_view const& covariance, stddev_iter, d_result, [] __device__(auto const covariance, auto const stddev) { - return covariance / thrust::get<0>(stddev) / thrust::get<1>(stddev); + return covariance / cuda::std::get<0>(stddev) / cuda::std::get<1>(stddev); }); result->set_null_count(covariance.null_count()); diff --git a/cpp/src/groupby/sort/group_merge_m2.cu b/cpp/src/groupby/sort/group_merge_m2.cu index a580c9dac9d..1bb92a82b58 100644 --- a/cpp/src/groupby/sort/group_merge_m2.cu +++ b/cpp/src/groupby/sort/group_merge_m2.cu @@ -32,7 +32,7 @@ #include #include #include -#include +#include namespace cudf { namespace groupby { @@ -90,7 +90,7 @@ struct merge_fn { // This case should never happen, because all groups are non-empty as the results of // aggregation. Here we just to make sure we cover this case. if (start_idx == end_idx) { - return thrust::make_tuple(size_type{0}, result_type{0}, result_type{0}, int8_t{0}); + return cuda::std::make_tuple(size_type{0}, result_type{0}, result_type{0}, int8_t{0}); } // If `(n = d_counts[idx]) > 0` then `d_means[idx] != null` and `d_M2s[idx] != null`. @@ -121,7 +121,7 @@ struct merge_fn { // zero), then the output is a null. auto const is_valid = int8_t{merge_vals.count > 0}; - return thrust::make_tuple(merge_vals.count, merge_vals.mean, merge_vals.M2, is_valid); + return cuda::std::make_tuple(merge_vals.count, merge_vals.mean, merge_vals.M2, is_valid); } }; @@ -157,10 +157,10 @@ std::unique_ptr group_merge_m2(column_view const& values, // Perform merging for all the aggregations. Their output (and their validity data) are written // out concurrently through an output zip iterator. - using iterator_tuple = thrust::tuple; + using iterator_tuple = cuda::std::tuple; using output_iterator = thrust::zip_iterator; auto const out_iter = - output_iterator{thrust::make_tuple(result_counts->mutable_view().template data(), + output_iterator{cuda::std::make_tuple(result_counts->mutable_view().template data(), result_means->mutable_view().template data(), result_M2s->mutable_view().template data(), validities.begin())}; diff --git a/cpp/src/groupby/sort/group_replace_nulls.cu b/cpp/src/groupby/sort/group_replace_nulls.cu index 49557164230..9d9611f5530 100644 --- a/cpp/src/groupby/sort/group_replace_nulls.cu +++ b/cpp/src/groupby/sort/group_replace_nulls.cu @@ -28,7 +28,7 @@ #include #include #include -#include +#include #include @@ -47,11 +47,11 @@ std::unique_ptr group_replace_nulls(cudf::column_view const& grouped_val auto device_in = cudf::column_device_view::create(grouped_value, stream); auto index = thrust::make_counting_iterator(0); auto valid_it = cudf::detail::make_validity_iterator(*device_in); - auto in_begin = thrust::make_zip_iterator(thrust::make_tuple(index, valid_it)); + auto in_begin = thrust::make_zip_iterator(cuda::std::make_tuple(index, valid_it)); rmm::device_uvector gather_map(size, stream); auto gm_begin = thrust::make_zip_iterator( - thrust::make_tuple(gather_map.begin(), thrust::make_discard_iterator())); + cuda::std::make_tuple(gather_map.begin(), thrust::make_discard_iterator())); auto func = cudf::detail::replace_policy_functor(); thrust::equal_to eq; diff --git a/cpp/src/io/comp/nvcomp_adapter.cu b/cpp/src/io/comp/nvcomp_adapter.cu index 794d452ebf2..60c4d084d76 100644 --- a/cpp/src/io/comp/nvcomp_adapter.cu +++ b/cpp/src/io/comp/nvcomp_adapter.cu @@ -21,7 +21,7 @@ #include #include -#include +#include namespace cudf::io::nvcomp { @@ -39,7 +39,7 @@ batched_args create_batched_nvcomp_args(device_span c auto ins_it = thrust::make_zip_iterator(input_data_ptrs.begin(), input_data_sizes.begin()); thrust::transform( rmm::exec_policy(stream), inputs.begin(), inputs.end(), ins_it, [] __device__(auto const& in) { - return thrust::make_tuple(in.data(), in.size()); + return cuda::std::make_tuple(in.data(), in.size()); }); // Prepare the output vectors @@ -49,7 +49,7 @@ batched_args create_batched_nvcomp_args(device_span c outputs.begin(), outputs.end(), outs_it, - [] __device__(auto const& out) { return thrust::make_tuple(out.data(), out.size()); }); + [] __device__(auto const& out) { return cuda::std::make_tuple(out.data(), out.size()); }); return {std::move(input_data_ptrs), std::move(input_data_sizes), diff --git a/cpp/src/io/comp/statistics.cu b/cpp/src/io/comp/statistics.cu index e0f7e1ec6dd..462622dcb1f 100644 --- a/cpp/src/io/comp/statistics.cu +++ b/cpp/src/io/comp/statistics.cu @@ -39,7 +39,7 @@ writer_compression_statistics collect_compression_statistics( auto input_size_with_status = [inputs, results, stream](compression_status status) { auto const zipped_begin = - thrust::make_zip_iterator(thrust::make_tuple(inputs.begin(), results.begin())); + thrust::make_zip_iterator(cuda::std::make_tuple(inputs.begin(), results.begin())); auto const zipped_end = zipped_begin + inputs.size(); return thrust::transform_reduce( @@ -47,7 +47,7 @@ writer_compression_statistics collect_compression_statistics( zipped_begin, zipped_end, [status] __device__(auto tup) { - return thrust::get<1>(tup).status == status ? thrust::get<0>(tup).size() : 0; + return cuda::std::get<1>(tup).status == status ? cuda::std::get<0>(tup).size() : 0; }, 0ul, thrust::plus()); diff --git a/cpp/src/io/json/json_column.cu b/cpp/src/io/json/json_column.cu index b1dc2c9dd7f..9a0b3a4f9e8 100644 --- a/cpp/src/io/json/json_column.cu +++ b/cpp/src/io/json/json_column.cu @@ -350,9 +350,9 @@ std::vector copy_strings_to_host(device_span input, thrust::make_zip_iterator(string_offsets.begin(), string_lengths.begin()), [] __device__(auto const& offsets) { // Note: first character for non-field columns - return thrust::make_tuple( - static_cast(thrust::get<0>(offsets)), - static_cast(thrust::get<1>(offsets) - thrust::get<0>(offsets))); + return cuda::std::make_tuple( + static_cast(cuda::std::get<0>(offsets)), + static_cast(cuda::std::get<1>(offsets) - cuda::std::get<0>(offsets))); }); cudf::io::parse_options_view options_view{}; @@ -546,7 +546,7 @@ void make_device_json_column(device_span input, auto h_range_col_id_it = thrust::make_zip_iterator(column_range_beg.begin(), unique_col_ids.begin()); std::sort(h_range_col_id_it, h_range_col_id_it + num_columns, [](auto const& a, auto const& b) { - return thrust::get<0>(a) < thrust::get<0>(b); + return cuda::std::get<0>(a) < cuda::std::get<0>(b); }); // use hash map because we may skip field name's col_ids @@ -667,7 +667,7 @@ void make_device_json_column(device_span input, // restore unique_col_ids order std::sort(h_range_col_id_it, h_range_col_id_it + num_columns, [](auto const& a, auto const& b) { - return thrust::get<1>(a) < thrust::get<1>(b); + return cuda::std::get<1>(a) < cuda::std::get<1>(b); }); // move columns data to device. std::vector columns_data(num_columns); diff --git a/cpp/src/io/json/json_tree.cu b/cpp/src/io/json/json_tree.cu index 148aeb5ec7a..78ae0a20671 100644 --- a/cpp/src/io/json/json_tree.cu +++ b/cpp/src/io/json/json_tree.cu @@ -84,7 +84,7 @@ struct node_ranges { device_span tokens; device_span token_indices; bool include_quote_char; - __device__ auto operator()(size_type i) -> thrust::tuple + __device__ auto operator()(size_type i) -> cuda::std::tuple { // Whether a token expects to be followed by its respective end-of-* token partner auto const is_begin_of_section = [] __device__(PdaTokenT const token) { @@ -129,7 +129,7 @@ struct node_ranges { range_end = get_token_index(tokens[i + 1], token_indices[i + 1]); } } - return thrust::make_tuple(range_begin, range_end); + return cuda::std::make_tuple(range_begin, range_end); } }; diff --git a/cpp/src/io/json/nested_json_gpu.cu b/cpp/src/io/json/nested_json_gpu.cu index 5eb3883dc64..1ec2fe20094 100644 --- a/cpp/src/io/json/nested_json_gpu.cu +++ b/cpp/src/io/json/nested_json_gpu.cu @@ -43,7 +43,7 @@ #include #include #include -#include +#include #include #include @@ -131,10 +131,10 @@ constexpr auto NUM_SYMBOL_GROUPS = static_cast(dfa_symbol_group_id::NU * @brief Function object to map (input_symbol,stack_context) tuples to a symbol group. */ struct SymbolPairToSymbolGroupId { - CUDF_HOST_DEVICE SymbolGroupT operator()(thrust::tuple symbol) const + CUDF_HOST_DEVICE SymbolGroupT operator()(cuda::std::tuple symbol) const { - auto const input_symbol = thrust::get<0>(symbol); - auto const stack_symbol = thrust::get<1>(symbol); + auto const input_symbol = cuda::std::get<0>(symbol); + auto const stack_symbol = cuda::std::get<1>(symbol); return static_cast( input_symbol == '\n' ? dfa_symbol_group_id::DELIMITER @@ -154,7 +154,7 @@ struct TransduceInputOp { SymbolT const read_symbol) const { if (state_id == static_cast(dfa_states::EXCESS)) { return '_'; } - return thrust::get<1>(read_symbol); + return cuda::std::get<1>(read_symbol); } template @@ -226,9 +226,9 @@ std::array, NUM_SYMBOL_GROUPS - 1> const symbol_groups{{ struct UnwrapTokenFromSymbolOp { template CUDF_HOST_DEVICE SymbolGroupT operator()(SymbolGroupLookupTableT const& sgid_lut, - thrust::tuple symbol) const + cuda::std::tuple symbol) const { - PdaTokenT const token_type = thrust::get<0>(symbol); + PdaTokenT const token_type = cuda::std::get<0>(symbol); return sgid_lut.lookup(token_type); } }; @@ -574,14 +574,14 @@ static __constant__ PdaSymbolGroupIdT tos_sg_to_pda_sgid[] = { struct PdaSymbolToSymbolGroupId { template __device__ __forceinline__ PdaSymbolGroupIdT - operator()(thrust::tuple symbol_pair) const + operator()(cuda::std::tuple symbol_pair) const { // The symbol read from the input - auto symbol = thrust::get<0>(symbol_pair); + auto symbol = cuda::std::get<0>(symbol_pair); // The stack symbol (i.e., what is on top of the stack at the time the input symbol was read) // I.e., whether we're reading in something within a struct, a list, or the JSON root - auto stack_symbol = thrust::get<1>(symbol_pair); + auto stack_symbol = cuda::std::get<1>(symbol_pair); // The stack symbol offset: '_' is the root group (0), '[' is the list group (1), '{' is the // struct group (2) diff --git a/cpp/src/io/json/write_json.cu b/cpp/src/io/json/write_json.cu index 84e0ac9e74d..31187f52945 100644 --- a/cpp/src/io/json/write_json.cu +++ b/cpp/src/io/json/write_json.cu @@ -700,8 +700,8 @@ struct column_to_strings_fn { i_col_begin + num_columns, std::back_inserter(str_column_vec), [this, &children_names](auto const& i_current_col) { - auto const i = thrust::get<0>(i_current_col); - auto const& current_col = thrust::get<1>(i_current_col); + auto const i = cuda::std::get<0>(i_current_col); + auto const& current_col = cuda::std::get<1>(i_current_col); // Struct needs children's column names if (current_col.type().id() == type_id::STRUCT) { return this->template operator()(current_col, diff --git a/cpp/src/io/orc/stripe_enc.cu b/cpp/src/io/orc/stripe_enc.cu index 516922219d1..b83fe9a006d 100644 --- a/cpp/src/io/orc/stripe_enc.cu +++ b/cpp/src/io/orc/stripe_enc.cu @@ -35,7 +35,7 @@ #include #include #include -#include +#include namespace cudf { namespace io { diff --git a/cpp/src/io/parquet/page_enc.cu b/cpp/src/io/parquet/page_enc.cu index 2f351edd2b9..8637fdc6258 100644 --- a/cpp/src/io/parquet/page_enc.cu +++ b/cpp/src/io/parquet/page_enc.cu @@ -41,7 +41,7 @@ #include #include #include -#include +#include #include diff --git a/cpp/src/io/parquet/page_hdr.cu b/cpp/src/io/parquet/page_hdr.cu index 888d9452612..624c82639af 100644 --- a/cpp/src/io/parquet/page_hdr.cu +++ b/cpp/src/io/parquet/page_hdr.cu @@ -20,7 +20,7 @@ #include -#include +#include #include @@ -241,10 +241,10 @@ struct FunctionSwitchImpl { static inline __device__ bool run(byte_stream_s* bs, int field_type, int const& field, - thrust::tuple& ops) + cuda::std::tuple& ops) { - if (field == thrust::get(ops).field) { - return thrust::get(ops)(bs, field_type); + if (field == cuda::std::get(ops).field) { + return cuda::std::get(ops)(bs, field_type); } else { return FunctionSwitchImpl::run(bs, field_type, field, ops); } @@ -257,10 +257,10 @@ struct FunctionSwitchImpl<0> { static inline __device__ bool run(byte_stream_s* bs, int field_type, int const& field, - thrust::tuple& ops) + cuda::std::tuple& ops) { - if (field == thrust::get<0>(ops).field) { - return thrust::get<0>(ops)(bs, field_type); + if (field == cuda::std::get<0>(ops).field) { + return cuda::std::get<0>(ops)(bs, field_type); } else { skip_struct_field(bs, field_type); return false; @@ -279,9 +279,9 @@ struct FunctionSwitchImpl<0> { * byte stream. Otherwise true is returned. */ template -inline __device__ bool parse_header(thrust::tuple& op, byte_stream_s* bs) +inline __device__ bool parse_header(cuda::std::tuple& op, byte_stream_s* bs) { - constexpr int index = thrust::tuple_size>::value - 1; + constexpr int index = cuda::std::tuple_size>::value - 1; int field = 0; while (true) { auto const current_byte = getb(bs); @@ -298,7 +298,7 @@ inline __device__ bool parse_header(thrust::tuple& op, byte_stream_ struct gpuParseDataPageHeader { __device__ bool operator()(byte_stream_s* bs) { - auto op = thrust::make_tuple(ParquetFieldInt32(1, bs->page.num_input_values), + auto op = cuda::std::make_tuple(ParquetFieldInt32(1, bs->page.num_input_values), ParquetFieldEnum(2, bs->page.encoding), ParquetFieldEnum(3, bs->page.definition_level_encoding), ParquetFieldEnum(4, bs->page.repetition_level_encoding)); @@ -309,7 +309,7 @@ struct gpuParseDataPageHeader { struct gpuParseDictionaryPageHeader { __device__ bool operator()(byte_stream_s* bs) { - auto op = thrust::make_tuple(ParquetFieldInt32(1, bs->page.num_input_values), + auto op = cuda::std::make_tuple(ParquetFieldInt32(1, bs->page.num_input_values), ParquetFieldEnum(2, bs->page.encoding)); return parse_header(op, bs); } @@ -318,7 +318,7 @@ struct gpuParseDictionaryPageHeader { struct gpuParseDataPageHeaderV2 { __device__ bool operator()(byte_stream_s* bs) { - auto op = thrust::make_tuple(ParquetFieldInt32(1, bs->page.num_input_values), + auto op = cuda::std::make_tuple(ParquetFieldInt32(1, bs->page.num_input_values), ParquetFieldInt32(2, bs->page.num_nulls), ParquetFieldInt32(3, bs->page.num_rows), ParquetFieldEnum(4, bs->page.encoding), @@ -331,7 +331,7 @@ struct gpuParseDataPageHeaderV2 { struct gpuParsePageHeader { __device__ bool operator()(byte_stream_s* bs) { - auto op = thrust::make_tuple(ParquetFieldEnum(1, bs->page_type), + auto op = cuda::std::make_tuple(ParquetFieldEnum(1, bs->page_type), ParquetFieldInt32(2, bs->page.uncompressed_page_size), ParquetFieldInt32(3, bs->page.compressed_page_size), ParquetFieldStruct(5), diff --git a/cpp/src/io/parquet/reader_impl_helpers.hpp b/cpp/src/io/parquet/reader_impl_helpers.hpp index 8d8ab8707be..ea569fed494 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.hpp +++ b/cpp/src/io/parquet/reader_impl_helpers.hpp @@ -245,7 +245,7 @@ class named_to_reference_converter : public ast::detail::expression_transformer thrust::counting_iterator(metadata.schema_info.size())), std::inserter(column_name_to_index, column_name_to_index.end()), [](auto const& name_index) { - return std::make_pair(thrust::get<0>(name_index).name, thrust::get<1>(name_index)); + return std::make_pair(cuda::std::get<0>(name_index).name, cuda::std::get<1>(name_index)); }); expr.value().get().accept(*this); diff --git a/cpp/src/io/text/bgzip_data_chunk_source.cu b/cpp/src/io/text/bgzip_data_chunk_source.cu index 77647c18b20..199080609ac 100644 --- a/cpp/src/io/text/bgzip_data_chunk_source.cu +++ b/cpp/src/io/text/bgzip_data_chunk_source.cu @@ -50,14 +50,14 @@ struct bgzip_nvcomp_transform_functor { uint8_t const* compressed_ptr; uint8_t* decompressed_ptr; - __device__ thrust::tuple, device_span> operator()( - thrust::tuple t) + __device__ cuda::std::tuple, device_span> operator()( + cuda::std::tuple t) { - auto const compressed_begin = thrust::get<0>(t); - auto const compressed_end = thrust::get<1>(t); - auto const decompressed_begin = thrust::get<2>(t); - auto const decompressed_end = thrust::get<3>(t); - return thrust::make_tuple(device_span{compressed_ptr + compressed_begin, + auto const compressed_begin = cuda::std::get<0>(t); + auto const compressed_end = cuda::std::get<1>(t); + auto const decompressed_begin = cuda::std::get<2>(t); + auto const decompressed_end = cuda::std::get<3>(t); + return cuda::std::make_tuple(device_span{compressed_ptr + compressed_begin, compressed_end - compressed_begin}, device_span{decompressed_ptr + decompressed_begin, decompressed_end - decompressed_begin}); diff --git a/cpp/src/io/utilities/data_casting.cu b/cpp/src/io/utilities/data_casting.cu index 8fd860d9492..67ca9eb8922 100644 --- a/cpp/src/io/utilities/data_casting.cu +++ b/cpp/src/io/utilities/data_casting.cu @@ -783,10 +783,10 @@ template struct to_string_view_pair { SymbolT const* data; to_string_view_pair(SymbolT const* _data) : data(_data) {} - __device__ auto operator()(thrust::tuple ip) + __device__ auto operator()(cuda::std::tuple ip) { - return thrust::pair{data + thrust::get<0>(ip), - static_cast(thrust::get<1>(ip))}; + return thrust::pair{data + cuda::std::get<0>(ip), + static_cast(cuda::std::get<1>(ip))}; } }; @@ -908,7 +908,7 @@ static std::unique_ptr parse_string(string_view_pair_it str_tuples, std::unique_ptr parse_data( const char* data, - thrust::zip_iterator> offset_length_begin, + thrust::zip_iterator> offset_length_begin, size_type col_size, data_type col_type, rmm::device_buffer&& null_mask, diff --git a/cpp/src/io/utilities/string_parsing.hpp b/cpp/src/io/utilities/string_parsing.hpp index 12fc0a5b2e7..e3c95540888 100644 --- a/cpp/src/io/utilities/string_parsing.hpp +++ b/cpp/src/io/utilities/string_parsing.hpp @@ -23,7 +23,7 @@ #include #include -#include +#include namespace cudf::io { namespace detail { @@ -45,7 +45,7 @@ namespace detail { cudf::data_type infer_data_type( cudf::io::json_inference_options_view const& options, device_span data, - thrust::zip_iterator> offset_length_begin, + thrust::zip_iterator> offset_length_begin, std::size_t const size, rmm::cuda_stream_view stream); } // namespace detail @@ -67,7 +67,7 @@ namespace json::detail { */ std::unique_ptr parse_data( const char* data, - thrust::zip_iterator> offset_length_begin, + thrust::zip_iterator> offset_length_begin, size_type col_size, data_type col_type, rmm::device_buffer&& null_mask, diff --git a/cpp/src/io/utilities/type_inference.cu b/cpp/src/io/utilities/type_inference.cu index b446ad41946..0254616c672 100644 --- a/cpp/src/io/utilities/type_inference.cu +++ b/cpp/src/io/utilities/type_inference.cu @@ -102,7 +102,7 @@ __device__ __inline__ bool is_like_float(std::size_t len, * @tparam BlockSize Number of threads in each block * @tparam OptionsView Type of inference options view * @tparam ColumnStringIter Iterator type whose `value_type` is a - * `thrust::tuple`, where `offset_t` and `length_t` are of integral type and + * `cuda::std::tuple`, where `offset_t` and `length_t` are of integral type and * `offset_t` needs to be convertible to `std::size_t`. * * @param[in] options View of inference options @@ -122,8 +122,8 @@ CUDF_KERNEL void infer_column_type_kernel(OptionsView options, for (auto idx = threadIdx.x + blockDim.x * blockIdx.x; idx < size; idx += gridDim.x * blockDim.x) { - auto const field_offset = thrust::get<0>(*(offset_length_begin + idx)); - auto const field_len = thrust::get<1>(*(offset_length_begin + idx)); + auto const field_offset = cuda::std::get<0>(*(offset_length_begin + idx)); + auto const field_len = cuda::std::get<1>(*(offset_length_begin + idx)); auto const field_begin = data.begin() + field_offset; if (cudf::detail::serialized_trie_contains( @@ -222,7 +222,7 @@ CUDF_KERNEL void infer_column_type_kernel(OptionsView options, * * @tparam OptionsView Type of inference options view * @tparam ColumnStringIter Iterator type whose `value_type` is a - * `thrust::tuple`, where `offset_t` and `length_t` are of integral type and + * `cuda::std::tuple`, where `offset_t` and `length_t` are of integral type and * `offset_t` needs to be convertible to `std::size_t`. * * @param options View of inference options @@ -255,7 +255,7 @@ cudf::io::column_type_histogram infer_column_type(OptionsView const& options, cudf::data_type infer_data_type( cudf::io::json_inference_options_view const& options, device_span data, - thrust::zip_iterator> offset_length_begin, + thrust::zip_iterator> offset_length_begin, std::size_t const size, rmm::cuda_stream_view stream) { diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 17616818a58..3f647fd7a6a 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -33,7 +33,7 @@ #include #include #include -#include +#include #include #include @@ -194,9 +194,9 @@ probe_join_hash_table( cudf::size_type const probe_table_num_rows = probe_table.num_rows(); auto const out1_zip_begin = thrust::make_zip_iterator( - thrust::make_tuple(thrust::make_discard_iterator(), left_indices->begin())); + cuda::std::make_tuple(thrust::make_discard_iterator(), left_indices->begin())); auto const out2_zip_begin = thrust::make_zip_iterator( - thrust::make_tuple(thrust::make_discard_iterator(), right_indices->begin())); + cuda::std::make_tuple(thrust::make_discard_iterator(), right_indices->begin())); auto const row_comparator = cudf::experimental::row::equality::two_table_comparator{preprocessed_probe, preprocessed_build}; @@ -296,9 +296,9 @@ std::size_t get_full_join_size( cudf::size_type const probe_table_num_rows = probe_table.num_rows(); auto const out1_zip_begin = thrust::make_zip_iterator( - thrust::make_tuple(thrust::make_discard_iterator(), left_indices->begin())); + cuda::std::make_tuple(thrust::make_discard_iterator(), left_indices->begin())); auto const out2_zip_begin = thrust::make_zip_iterator( - thrust::make_tuple(thrust::make_discard_iterator(), right_indices->begin())); + cuda::std::make_tuple(thrust::make_discard_iterator(), right_indices->begin())); auto const row_comparator = cudf::experimental::row::equality::two_table_comparator{preprocessed_probe, preprocessed_build}; diff --git a/cpp/src/json/json_path.cu b/cpp/src/json/json_path.cu index 2be5798098d..a48423ed7eb 100644 --- a/cpp/src/json/json_path.cu +++ b/cpp/src/json/json_path.cu @@ -41,7 +41,7 @@ #include #include #include -#include +#include namespace cudf { namespace detail { diff --git a/cpp/src/lists/dremel.cu b/cpp/src/lists/dremel.cu index 7a460d3dfab..51f708ba510 100644 --- a/cpp/src/lists/dremel.cu +++ b/cpp/src/lists/dremel.cu @@ -318,13 +318,13 @@ dremel_data get_encoding(column_view h_col, // Zip the input and output value iterators so that merge operation is done only once auto input_parent_zip_it = - thrust::make_zip_iterator(thrust::make_tuple(input_parent_rep_it, input_parent_def_it)); + thrust::make_zip_iterator(cuda::std::make_tuple(input_parent_rep_it, input_parent_def_it)); auto input_child_zip_it = - thrust::make_zip_iterator(thrust::make_tuple(input_child_rep_it, input_child_def_it)); + thrust::make_zip_iterator(cuda::std::make_tuple(input_child_rep_it, input_child_def_it)); auto output_zip_it = - thrust::make_zip_iterator(thrust::make_tuple(rep_level.begin(), def_level.begin())); + thrust::make_zip_iterator(cuda::std::make_tuple(rep_level.begin(), def_level.begin())); auto ends = thrust::merge_by_key(rmm::exec_policy(stream), empties.begin(), @@ -405,13 +405,13 @@ dremel_data get_encoding(column_view h_col, // Zip the input and output value iterators so that merge operation is done only once auto input_parent_zip_it = - thrust::make_zip_iterator(thrust::make_tuple(input_parent_rep_it, input_parent_def_it)); + thrust::make_zip_iterator(cuda::std::make_tuple(input_parent_rep_it, input_parent_def_it)); auto input_child_zip_it = - thrust::make_zip_iterator(thrust::make_tuple(temp_rep_vals.begin(), temp_def_vals.begin())); + thrust::make_zip_iterator(cuda::std::make_tuple(temp_rep_vals.begin(), temp_def_vals.begin())); auto output_zip_it = - thrust::make_zip_iterator(thrust::make_tuple(rep_level.begin(), def_level.begin())); + thrust::make_zip_iterator(cuda::std::make_tuple(rep_level.begin(), def_level.begin())); auto ends = thrust::merge_by_key(rmm::exec_policy(stream), transformed_empties, diff --git a/cpp/src/merge/merge.cu b/cpp/src/merge/merge.cu index 073a2a6b97e..2d9bafc54c7 100644 --- a/cpp/src/merge/merge.cu +++ b/cpp/src/merge/merge.cu @@ -48,7 +48,7 @@ #include #include #include -#include +#include #include @@ -106,8 +106,8 @@ using index_type = detail::index_type; * * Merges the bits from two column_device_views into the destination validity buffer * according to `merged_indices` map such that bit `i` in `out_validity` - * will be equal to bit `thrust::get<1>(merged_indices[i])` from `left_dcol` - * if `thrust::get<0>(merged_indices[i])` equals `side::LEFT`; otherwise, + * will be equal to bit `cuda::std::get<1>(merged_indices[i])` from `left_dcol` + * if `cuda::std::get<0>(merged_indices[i])` equals `side::LEFT`; otherwise, * from `right_dcol`. * * `left_dcol` and `right_dcol` must not overlap. diff --git a/cpp/src/partitioning/round_robin.cu b/cpp/src/partitioning/round_robin.cu index c615f08ff12..c8b0f9b290f 100644 --- a/cpp/src/partitioning/round_robin.cu +++ b/cpp/src/partitioning/round_robin.cu @@ -38,7 +38,7 @@ #include #include #include -#include +#include #include diff --git a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu index fc56d17d73b..7450542af75 100644 --- a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu +++ b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu @@ -50,7 +50,7 @@ #include #include #include -#include +#include #include @@ -64,7 +64,7 @@ namespace { // values. {mean, weight} // NOTE: Using a tuple here instead of a struct to take advantage of // thrust zip iterators for output. -using centroid = thrust::tuple; +using centroid = cuda::std::tuple; // make a centroid from a scalar with a weight of 1. template @@ -104,16 +104,16 @@ struct make_weighted_centroid { struct merge_centroids { centroid operator() __device__(centroid const& lhs, centroid const& rhs) const { - bool const lhs_valid = thrust::get<2>(lhs); - bool const rhs_valid = thrust::get<2>(rhs); + bool const lhs_valid = cuda::std::get<2>(lhs); + bool const rhs_valid = cuda::std::get<2>(rhs); if (!lhs_valid && !rhs_valid) { return {0, 0, false}; } if (!lhs_valid) { return rhs; } if (!rhs_valid) { return lhs; } - double const lhs_mean = thrust::get<0>(lhs); - double const rhs_mean = thrust::get<0>(rhs); - double const lhs_weight = thrust::get<1>(lhs); - double const rhs_weight = thrust::get<1>(rhs); + double const lhs_mean = cuda::std::get<0>(lhs); + double const rhs_mean = cuda::std::get<0>(rhs); + double const lhs_weight = cuda::std::get<1>(lhs); + double const rhs_weight = cuda::std::get<1>(rhs); double const new_weight = lhs_weight + rhs_weight; return {(lhs_mean * lhs_weight + rhs_mean * rhs_weight) / new_weight, new_weight, true}; } @@ -261,7 +261,7 @@ struct scalar_group_info_grouped { size_type const* group_valid_counts; size_type const* group_offsets; - __device__ thrust::tuple operator()(size_type group_index) const + __device__ cuda::std::tuple operator()(size_type group_index) const { return {static_cast(group_valid_counts[group_index]), group_offsets[group_index + 1] - group_offsets[group_index], @@ -274,7 +274,7 @@ struct scalar_group_info { double const total_weight; size_type const size; - __device__ thrust::tuple operator()(size_type) const + __device__ cuda::std::tuple operator()(size_type) const { return {total_weight, size, 0}; } @@ -287,7 +287,7 @@ struct centroid_group_info { GroupOffsetsIter outer_offsets; size_type const* inner_offsets; - __device__ thrust::tuple operator()(size_type group_index) const + __device__ cuda::std::tuple operator()(size_type group_index) const { // if there's no weights in this group of digests at all, return 0. auto const group_start = inner_offsets[outer_offsets[group_index]]; @@ -295,26 +295,26 @@ struct centroid_group_info { auto const num_weights = group_end - group_start; auto const last_weight_index = group_end - 1; return num_weights == 0 - ? thrust::tuple{0, num_weights, group_start} - : thrust::tuple{ + ? cuda::std::tuple{0, num_weights, group_start} + : cuda::std::tuple{ cumulative_weights[last_weight_index], num_weights, group_start}; } }; struct tdigest_min { - __device__ double operator()(thrust::tuple const& t) const + __device__ double operator()(cuda::std::tuple const& t) const { - auto const min = thrust::get<0>(t); - auto const size = thrust::get<1>(t); + auto const min = cuda::std::get<0>(t); + auto const size = cuda::std::get<1>(t); return size > 0 ? min : std::numeric_limits::max(); } }; struct tdigest_max { - __device__ double operator()(thrust::tuple const& t) const + __device__ double operator()(cuda::std::tuple const& t) const { - auto const max = thrust::get<0>(t); - auto const size = thrust::get<1>(t); + auto const max = cuda::std::get<0>(t); + auto const size = cuda::std::get<1>(t); return size > 0 ? max : std::numeric_limits::lowest(); } }; @@ -772,7 +772,7 @@ std::unique_ptr compute_tdigests(int delta, cudf::mutable_column_view weight_col(*centroid_weights); // reduce the centroids into the clusters - auto output = thrust::make_zip_iterator(thrust::make_tuple( + auto output = thrust::make_zip_iterator(cuda::std::make_tuple( mean_col.begin(), weight_col.begin(), thrust::make_discard_iterator())); auto const num_values = std::distance(centroids_begin, centroids_end); @@ -804,14 +804,14 @@ struct get_scalar_minmax_grouped { device_span group_offsets; size_type const* group_valid_counts; - __device__ thrust::tuple operator()(size_type group_index) + __device__ cuda::std::tuple operator()(size_type group_index) { auto const valid_count = group_valid_counts[group_index]; return valid_count > 0 - ? thrust::make_tuple( + ? cuda::std::make_tuple( static_cast(col.element(group_offsets[group_index])), static_cast(col.element(group_offsets[group_index] + valid_count - 1))) - : thrust::make_tuple(0.0, 0.0); + : cuda::std::make_tuple(0.0, 0.0); } }; @@ -821,12 +821,12 @@ struct get_scalar_minmax { column_device_view const col; size_type const valid_count; - __device__ thrust::tuple operator()(size_type) + __device__ cuda::std::tuple operator()(size_type) { return valid_count > 0 - ? thrust::make_tuple(static_cast(col.element(0)), + ? cuda::std::make_tuple(static_cast(col.element(0)), static_cast(col.element(valid_count - 1))) - : thrust::make_tuple(0.0, 0.0); + : cuda::std::make_tuple(0.0, 0.0); } }; @@ -866,7 +866,7 @@ struct typed_group_tdigest { rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + num_groups, - thrust::make_zip_iterator(thrust::make_tuple(min_col->mutable_view().begin(), + thrust::make_zip_iterator(cuda::std::make_tuple(min_col->mutable_view().begin(), max_col->mutable_view().begin())), get_scalar_minmax_grouped{*d_col, group_offsets, group_valid_counts.begin()}); @@ -946,7 +946,7 @@ struct typed_reduce_tdigest { rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + 1, - thrust::make_zip_iterator(thrust::make_tuple(min_col->mutable_view().begin(), + thrust::make_zip_iterator(cuda::std::make_tuple(min_col->mutable_view().begin(), max_col->mutable_view().begin())), get_scalar_minmax{*d_col, valid_count}); @@ -1087,7 +1087,7 @@ std::unique_ptr merge_tdigests(tdigest_column_view const& tdv, auto merged_min_col = cudf::make_numeric_column( data_type{type_id::FLOAT64}, num_groups, mask_state::UNALLOCATED, stream, mr); auto min_iter = - thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple( + thrust::make_transform_iterator(thrust::make_zip_iterator(cuda::std::make_tuple( tdv.min_begin(), cudf::tdigest::detail::size_begin(tdv))), tdigest_min{}); thrust::reduce_by_key(rmm::exec_policy(stream), @@ -1102,7 +1102,7 @@ std::unique_ptr merge_tdigests(tdigest_column_view const& tdv, auto merged_max_col = cudf::make_numeric_column( data_type{type_id::FLOAT64}, num_groups, mask_state::UNALLOCATED, stream, mr); auto max_iter = - thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple( + thrust::make_transform_iterator(thrust::make_zip_iterator(cuda::std::make_tuple( tdv.max_begin(), cudf::tdigest::detail::size_begin(tdv))), tdigest_max{}); thrust::reduce_by_key(rmm::exec_policy(stream), diff --git a/cpp/src/reductions/histogram.cu b/cpp/src/reductions/histogram.cu index 42ef266a684..042706b9b4d 100644 --- a/cpp/src/reductions/histogram.cu +++ b/cpp/src/reductions/histogram.cu @@ -23,7 +23,7 @@ #include #include -#include +#include #include #include @@ -96,7 +96,7 @@ struct is_not_zero { template __device__ bool operator()(Pair const input) const { - return thrust::get<1>(input) != 0; + return cuda::std::get<1>(input) != 0; } }; @@ -223,8 +223,8 @@ compute_row_frequencies(table_view const& input, rmm::mr::get_current_device_resource()); auto const input_it = thrust::make_zip_iterator( - thrust::make_tuple(thrust::make_counting_iterator(0), reduction_results.begin())); - auto const output_it = thrust::make_zip_iterator(thrust::make_tuple( + cuda::std::make_tuple(thrust::make_counting_iterator(0), reduction_results.begin())); + auto const output_it = thrust::make_zip_iterator(cuda::std::make_tuple( distinct_indices->begin(), distinct_counts->mutable_view().begin())); // Reduction results above are either group sizes of equal rows, or `0`. diff --git a/cpp/src/replace/clamp.cu b/cpp/src/replace/clamp.cu index 23c792ddcae..aabc1fc6656 100644 --- a/cpp/src/replace/clamp.cu +++ b/cpp/src/replace/clamp.cu @@ -42,7 +42,7 @@ #include #include #include -#include +#include #include @@ -132,17 +132,17 @@ std::enable_if_t(), std::unique_ptr> clamp cudf::mutable_column_device_view::create(output->mutable_view(), stream); auto input_device_view = cudf::column_device_view::create(input, stream); auto scalar_zip_itr = - thrust::make_zip_iterator(thrust::make_tuple(lo_itr, lo_replace_itr, hi_itr, hi_replace_itr)); + thrust::make_zip_iterator(cuda::std::make_tuple(lo_itr, lo_replace_itr, hi_itr, hi_replace_itr)); auto trans = cuda::proclaim_return_type([] __device__(auto element_optional, auto scalar_tuple) { if (element_optional.has_value()) { - auto lo_optional = thrust::get<0>(scalar_tuple); - auto hi_optional = thrust::get<2>(scalar_tuple); + auto lo_optional = cuda::std::get<0>(scalar_tuple); + auto hi_optional = cuda::std::get<2>(scalar_tuple); if (lo_optional.has_value() and (*element_optional < *lo_optional)) { - return *(thrust::get<1>(scalar_tuple)); + return *(cuda::std::get<1>(scalar_tuple)); } else if (hi_optional.has_value() and (*element_optional > *hi_optional)) { - return *(thrust::get<3>(scalar_tuple)); + return *(cuda::std::get<3>(scalar_tuple)); } } diff --git a/cpp/src/replace/nulls.cu b/cpp/src/replace/nulls.cu index 8ea229368cc..3a9aaf6aded 100644 --- a/cpp/src/replace/nulls.cu +++ b/cpp/src/replace/nulls.cu @@ -50,7 +50,7 @@ #include #include #include -#include +#include namespace { // anonymous @@ -367,11 +367,11 @@ std::unique_ptr replace_nulls_policy_impl(cudf::column_view const& auto device_in = cudf::column_device_view::create(input, stream); auto index = thrust::make_counting_iterator(0); auto valid_it = cudf::detail::make_validity_iterator(*device_in); - auto in_begin = thrust::make_zip_iterator(thrust::make_tuple(index, valid_it)); + auto in_begin = thrust::make_zip_iterator(cuda::std::make_tuple(index, valid_it)); rmm::device_uvector gather_map(input.size(), stream); auto gm_begin = thrust::make_zip_iterator( - thrust::make_tuple(gather_map.begin(), thrust::make_discard_iterator())); + cuda::std::make_tuple(gather_map.begin(), thrust::make_discard_iterator())); auto func = cudf::detail::replace_policy_functor(); if (replace_policy == cudf::replace_policy::PRECEDING) { diff --git a/cpp/src/replace/replace.cu b/cpp/src/replace/replace.cu index 184c30246c7..7b30926855c 100644 --- a/cpp/src/replace/replace.cu +++ b/cpp/src/replace/replace.cu @@ -57,7 +57,7 @@ #include #include #include -#include +#include namespace { // anonymous diff --git a/cpp/src/sort/rank.cu b/cpp/src/sort/rank.cu index 9cf07f065d2..e50d10fd53f 100644 --- a/cpp/src/sort/rank.cu +++ b/cpp/src/sort/rank.cu @@ -42,7 +42,7 @@ #include #include #include -#include +#include #include #include @@ -256,8 +256,8 @@ void rank_average(cudf::device_span group_keys, rank_count1.second + rank_count2.second}; }), cuda::proclaim_return_type([] __device__(MinCount minrank_count) { // min+(count-1)/2 - return static_cast(thrust::get<0>(minrank_count)) + - (static_cast(thrust::get<1>(minrank_count)) - 1) / 2.0; + return static_cast(cuda::std::get<0>(minrank_count)) + + (static_cast(cuda::std::get<1>(minrank_count)) - 1) / 2.0; }), stream); } diff --git a/cpp/src/strings/split/split_re.cu b/cpp/src/strings/split/split_re.cu index d8385549840..6b6f888a03f 100644 --- a/cpp/src/strings/split/split_re.cu +++ b/cpp/src/strings/split/split_re.cu @@ -78,7 +78,7 @@ struct token_reader_fn { auto const match = prog.find(prog_idx, d_str, itr); if (!match) { break; } - auto const start_pos = thrust::get<0>(match_positions_to_bytes(*match, d_str, last_pos)); + auto const start_pos = cuda::std::get<0>(match_positions_to_bytes(*match, d_str, last_pos)); // get the token (characters just before this match) auto const token = string_index_pair{d_str.data() + last_pos.byte_offset(), diff --git a/cpp/src/text/minhash.cu b/cpp/src/text/minhash.cu index dcb59166cec..5ae48c6238d 100644 --- a/cpp/src/text/minhash.cu +++ b/cpp/src/text/minhash.cu @@ -112,7 +112,7 @@ CUDF_KERNEL void minhash_kernel(cudf::column_device_view const d_strings, } else { // This code path assumes the use of MurmurHash3_x64_128 which produces 2 uint64 values // but only uses the first uint64 value as requested by the LLM team. - auto const hvalue = thrust::get<0>(hasher(hash_str)); + auto const hvalue = cuda::std::get<0>(hasher(hash_str)); cuda::atomic_ref ref{*(d_output + seed_idx)}; ref.fetch_min(hvalue, cuda::std::memory_order_relaxed); } diff --git a/cpp/tests/io/json_type_cast_test.cu b/cpp/tests/io/json_type_cast_test.cu index 6923b7be42d..73676fd24ef 100644 --- a/cpp/tests/io/json_type_cast_test.cu +++ b/cpp/tests/io/json_type_cast_test.cu @@ -44,9 +44,9 @@ struct JSONTypeCastTest : public cudf::test::BaseFixture {}; namespace { struct offsets_to_length { - __device__ cudf::size_type operator()(thrust::tuple const& p) + __device__ cudf::size_type operator()(cuda::std::tuple const& p) { - return thrust::get<1>(p) - thrust::get<0>(p); + return cuda::std::get<1>(p) - cuda::std::get<0>(p); } }; @@ -55,7 +55,7 @@ auto string_offset_to_length(cudf::strings_column_view const& column, rmm::cuda_ { auto offsets_begin = column.offsets_begin(); auto offsets_pair = - thrust::make_zip_iterator(thrust::make_tuple(offsets_begin, thrust::next(offsets_begin))); + thrust::make_zip_iterator(cuda::std::make_tuple(offsets_begin, thrust::next(offsets_begin))); rmm::device_uvector svs_length(column.size(), stream); thrust::transform(rmm::exec_policy(cudf::get_default_stream()), offsets_pair, @@ -96,7 +96,7 @@ TEST_F(JSONTypeCastTest, String) auto str_col = cudf::io::json::detail::parse_data( column.chars_begin(stream), - thrust::make_zip_iterator(thrust::make_tuple(column.offsets_begin(), svs_length.begin())), + thrust::make_zip_iterator(cuda::std::make_tuple(column.offsets_begin(), svs_length.begin())), column.size(), type, std::move(null_mask), @@ -129,7 +129,7 @@ TEST_F(JSONTypeCastTest, Int) auto col = cudf::io::json::detail::parse_data( column.chars_begin(stream), - thrust::make_zip_iterator(thrust::make_tuple(column.offsets_begin(), svs_length.begin())), + thrust::make_zip_iterator(cuda::std::make_tuple(column.offsets_begin(), svs_length.begin())), column.size(), type, std::move(null_mask), @@ -169,7 +169,7 @@ TEST_F(JSONTypeCastTest, StringEscapes) auto col = cudf::io::json::detail::parse_data( column.chars_begin(stream), - thrust::make_zip_iterator(thrust::make_tuple(column.offsets_begin(), svs_length.begin())), + thrust::make_zip_iterator(cuda::std::make_tuple(column.offsets_begin(), svs_length.begin())), column.size(), type, std::move(null_mask), @@ -238,7 +238,7 @@ TEST_F(JSONTypeCastTest, ErrorNulls) auto str_col = cudf::io::json::detail::parse_data( column.chars_begin(stream), - thrust::make_zip_iterator(thrust::make_tuple(column.offsets_begin(), svs_length.begin())), + thrust::make_zip_iterator(cuda::std::make_tuple(column.offsets_begin(), svs_length.begin())), column.size(), type, std::move(null_mask), diff --git a/cpp/tests/io/nested_json_test.cpp b/cpp/tests/io/nested_json_test.cpp index 070ac5ce870..3ee8626c93d 100644 --- a/cpp/tests/io/nested_json_test.cpp +++ b/cpp/tests/io/nested_json_test.cpp @@ -750,7 +750,7 @@ TEST_F(JsonTest, PostProcessTokenStream) // Golden token stream sample using token_t = cuio_json::token_t; using token_index_t = cuio_json::SymbolOffsetT; - using tuple_t = thrust::tuple; + using tuple_t = cuda::std::tuple; std::vector const input = {// Line 0 (invalid) {0, token_t::LineEnd}, @@ -861,9 +861,9 @@ TEST_F(JsonTest, PostProcessTokenStream) for (std::size_t i = 0; i < filtered_tokens.size(); i++) { // Ensure the index the tokens are pointing to do match - EXPECT_EQ(thrust::get<0>(expected_output[i]), filtered_indices[i]) << "Mismatch at #" << i; + EXPECT_EQ(cuda::std::get<0>(expected_output[i]), filtered_indices[i]) << "Mismatch at #" << i; // Ensure the token category is correct - EXPECT_EQ(thrust::get<1>(expected_output[i]), filtered_tokens[i]) << "Mismatch at #" << i; + EXPECT_EQ(cuda::std::get<1>(expected_output[i]), filtered_tokens[i]) << "Mismatch at #" << i; } } diff --git a/cpp/tests/io/type_inference_test.cu b/cpp/tests/io/type_inference_test.cu index 3bb15a59aa3..3b30ea1d965 100644 --- a/cpp/tests/io/type_inference_test.cu +++ b/cpp/tests/io/type_inference_test.cu @@ -26,7 +26,7 @@ #include #include -#include +#include #include #include @@ -59,7 +59,7 @@ TEST_F(TypeInference, Basic) string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); auto d_col_strings = - thrust::make_zip_iterator(thrust::make_tuple(d_string_offset.begin(), d_string_length.begin())); + thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -92,7 +92,7 @@ TEST_F(TypeInference, Null) string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); auto d_col_strings = - thrust::make_zip_iterator(thrust::make_tuple(d_string_offset.begin(), d_string_length.begin())); + thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -125,7 +125,7 @@ TEST_F(TypeInference, AllNull) string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); auto d_col_strings = - thrust::make_zip_iterator(thrust::make_tuple(d_string_offset.begin(), d_string_length.begin())); + thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -158,7 +158,7 @@ TEST_F(TypeInference, String) string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); auto d_col_strings = - thrust::make_zip_iterator(thrust::make_tuple(d_string_offset.begin(), d_string_length.begin())); + thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -191,7 +191,7 @@ TEST_F(TypeInference, Bool) string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); auto d_col_strings = - thrust::make_zip_iterator(thrust::make_tuple(d_string_offset.begin(), d_string_length.begin())); + thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -224,7 +224,7 @@ TEST_F(TypeInference, Timestamp) string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); auto d_col_strings = - thrust::make_zip_iterator(thrust::make_tuple(d_string_offset.begin(), d_string_length.begin())); + thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -258,7 +258,7 @@ TEST_F(TypeInference, InvalidInput) string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); auto d_col_strings = - thrust::make_zip_iterator(thrust::make_tuple(d_string_offset.begin(), d_string_length.begin())); + thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), diff --git a/cpp/tests/reductions/scan_tests.cpp b/cpp/tests/reductions/scan_tests.cpp index 161b1ee61ac..45df895dbbf 100644 --- a/cpp/tests/reductions/scan_tests.cpp +++ b/cpp/tests/reductions/scan_tests.cpp @@ -27,7 +27,7 @@ #include #include -#include +#include #include #include @@ -168,14 +168,14 @@ struct ScanTest : public BaseScanTest { bool const nullable = (b.size() > 0); auto masked_value = [identity](auto const& z) { - return thrust::get<1>(z) ? thrust::get<0>(z) : identity; + return cuda::std::get<1>(z) ? cuda::std::get<0>(z) : identity; }; if (inclusive == scan_type::INCLUSIVE) { if (nullable) { std::transform_inclusive_scan( - thrust::make_zip_iterator(thrust::make_tuple(v.begin(), b.begin())), - thrust::make_zip_iterator(thrust::make_tuple(v.end(), b.end())), + thrust::make_zip_iterator(cuda::std::make_tuple(v.begin(), b.begin())), + thrust::make_zip_iterator(cuda::std::make_tuple(v.end(), b.end())), expected.begin(), op, masked_value); @@ -189,8 +189,8 @@ struct ScanTest : public BaseScanTest { } else { if (nullable) { std::transform_exclusive_scan( - thrust::make_zip_iterator(thrust::make_tuple(v.begin(), b.begin())), - thrust::make_zip_iterator(thrust::make_tuple(v.end(), b.end())), + thrust::make_zip_iterator(cuda::std::make_tuple(v.begin(), b.begin())), + thrust::make_zip_iterator(cuda::std::make_tuple(v.end(), b.end())), expected.begin(), identity, op, diff --git a/cpp/tests/utilities/tdigest_utilities.cu b/cpp/tests/utilities/tdigest_utilities.cu index 9294aa0f681..193ec4ce7c5 100644 --- a/cpp/tests/utilities/tdigest_utilities.cu +++ b/cpp/tests/utilities/tdigest_utilities.cu @@ -29,7 +29,7 @@ #include #include #include -#include +#include // for use with groupby and reduction aggregation tests. @@ -58,9 +58,9 @@ void tdigest_sample_compare(cudf::tdigest::tdigest_column_view const& tdv, { auto iter = thrust::make_counting_iterator(0); std::for_each_n(iter, h_expected.size(), [&](size_type const index) { - h_expected_src[index] = thrust::get<0>(h_expected[index]); - h_expected_mean[index] = thrust::get<1>(h_expected[index]); - h_expected_weight[index] = thrust::get<2>(h_expected[index]); + h_expected_src[index] = cuda::std::get<0>(h_expected[index]); + h_expected_mean[index] = cuda::std::get<1>(h_expected[index]); + h_expected_weight[index] = cuda::std::get<2>(h_expected[index]); }); } From 6894f4099df857d8b64016d111904eb92af1a3a5 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Wed, 21 Feb 2024 11:13:01 -0600 Subject: [PATCH 2/2] Style. --- cpp/benchmarks/common/generate_input.cu | 2 +- cpp/benchmarks/string/url_decode.cu | 2 +- cpp/include/cudf/detail/merge.hpp | 2 +- cpp/include/cudf/detail/null_mask.cuh | 2 +- cpp/include/cudf/detail/replace/nulls.cuh | 4 +-- cpp/include/cudf/strings/detail/merge.cuh | 2 +- .../detail/strings_column_factories.cuh | 2 +- cpp/include/cudf_test/tdigest_utilities.cuh | 4 +-- cpp/src/copying/contiguous_split.cu | 2 +- cpp/src/groupby/sort/group_correlation.cu | 4 +-- cpp/src/groupby/sort/group_merge_m2.cu | 10 +++--- cpp/src/groupby/sort/group_replace_nulls.cu | 4 +-- cpp/src/io/comp/nvcomp_adapter.cu | 4 +-- cpp/src/io/comp/statistics.cu | 2 +- cpp/src/io/json/json_column.cu | 21 ++++++------ cpp/src/io/json/nested_json_gpu.cu | 4 +-- cpp/src/io/orc/stripe_enc.cu | 2 +- cpp/src/io/parquet/page_enc.cu | 2 +- cpp/src/io/parquet/page_hdr.cu | 34 ++++++++++--------- cpp/src/io/parquet/reader_impl_helpers.hpp | 2 +- cpp/src/io/text/bgzip_data_chunk_source.cu | 8 ++--- cpp/src/io/utilities/string_parsing.hpp | 4 +-- cpp/src/join/hash_join.cu | 2 +- cpp/src/json/json_path.cu | 2 +- cpp/src/lists/dremel.cu | 6 ++-- cpp/src/merge/merge.cu | 2 +- cpp/src/partitioning/round_robin.cu | 4 +-- .../quantiles/tdigest/tdigest_aggregation.cu | 8 ++--- cpp/src/reductions/histogram.cu | 2 +- cpp/src/replace/clamp.cu | 6 ++-- cpp/src/replace/nulls.cu | 2 +- cpp/src/replace/replace.cu | 2 +- cpp/src/sort/rank.cu | 4 +-- cpp/tests/io/type_inference_test.cu | 30 ++++++++-------- cpp/tests/reductions/scan_tests.cpp | 4 +-- cpp/tests/utilities/tdigest_utilities.cu | 4 +-- 36 files changed, 102 insertions(+), 99 deletions(-) diff --git a/cpp/benchmarks/common/generate_input.cu b/cpp/benchmarks/common/generate_input.cu index 567cabf4aa0..3fb455c0bbd 100644 --- a/cpp/benchmarks/common/generate_input.cu +++ b/cpp/benchmarks/common/generate_input.cu @@ -33,6 +33,7 @@ #include #include +#include #include #include #include @@ -51,7 +52,6 @@ #include #include #include -#include #include diff --git a/cpp/benchmarks/string/url_decode.cu b/cpp/benchmarks/string/url_decode.cu index 36d1c7e8ace..c1471a69643 100644 --- a/cpp/benchmarks/string/url_decode.cu +++ b/cpp/benchmarks/string/url_decode.cu @@ -26,12 +26,12 @@ #include #include +#include #include #include #include #include #include -#include struct url_string_generator { char* chars; diff --git a/cpp/include/cudf/detail/merge.hpp b/cpp/include/cudf/detail/merge.hpp index 27f94d085d1..e391d2f866b 100644 --- a/cpp/include/cudf/detail/merge.hpp +++ b/cpp/include/cudf/detail/merge.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2023, NVIDIA CORPORATION. + * Copyright (c) 2018-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/include/cudf/detail/null_mask.cuh b/cpp/include/cudf/detail/null_mask.cuh index f62679abac5..354d54e3a89 100644 --- a/cpp/include/cudf/detail/null_mask.cuh +++ b/cpp/include/cudf/detail/null_mask.cuh @@ -30,12 +30,12 @@ #include #include +#include #include #include #include #include #include -#include #include diff --git a/cpp/include/cudf/detail/replace/nulls.cuh b/cpp/include/cudf/detail/replace/nulls.cuh index f5061ee5a9a..f74f50a8cf7 100644 --- a/cpp/include/cudf/detail/replace/nulls.cuh +++ b/cpp/include/cudf/detail/replace/nulls.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -17,8 +17,8 @@ #include -#include #include +#include namespace cudf { namespace detail { diff --git a/cpp/include/cudf/strings/detail/merge.cuh b/cpp/include/cudf/strings/detail/merge.cuh index bc40c5472e7..31ae071911c 100644 --- a/cpp/include/cudf/strings/detail/merge.cuh +++ b/cpp/include/cudf/strings/detail/merge.cuh @@ -27,10 +27,10 @@ #include #include +#include #include #include #include -#include #include diff --git a/cpp/include/cudf/strings/detail/strings_column_factories.cuh b/cpp/include/cudf/strings/detail/strings_column_factories.cuh index 8884a5fca32..5185fb645b6 100644 --- a/cpp/include/cudf/strings/detail/strings_column_factories.cuh +++ b/cpp/include/cudf/strings/detail/strings_column_factories.cuh @@ -26,6 +26,7 @@ #include #include +#include #include #include #include @@ -35,7 +36,6 @@ #include #include #include -#include #include diff --git a/cpp/include/cudf_test/tdigest_utilities.cuh b/cpp/include/cudf_test/tdigest_utilities.cuh index 1b68262c38d..31cd142790e 100644 --- a/cpp/include/cudf_test/tdigest_utilities.cuh +++ b/cpp/include/cudf_test/tdigest_utilities.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -27,13 +27,13 @@ #include #include +#include #include #include #include #include #include #include -#include // for use with groupby and reduction aggregation tests. diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index 3bf105440c8..6c2513f29d4 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -32,6 +32,7 @@ #include #include +#include #include #include #include @@ -43,7 +44,6 @@ #include #include #include -#include #include diff --git a/cpp/src/groupby/sort/group_correlation.cu b/cpp/src/groupby/sort/group_correlation.cu index 0b1691dcc27..96b9b8f597b 100644 --- a/cpp/src/groupby/sort/group_correlation.cu +++ b/cpp/src/groupby/sort/group_correlation.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -27,13 +27,13 @@ #include #include +#include #include #include #include #include #include #include -#include #include diff --git a/cpp/src/groupby/sort/group_merge_m2.cu b/cpp/src/groupby/sort/group_merge_m2.cu index 1bb92a82b58..2e8206427f5 100644 --- a/cpp/src/groupby/sort/group_merge_m2.cu +++ b/cpp/src/groupby/sort/group_merge_m2.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -26,13 +26,13 @@ #include #include +#include #include #include #include #include #include #include -#include namespace cudf { namespace groupby { @@ -161,9 +161,9 @@ std::unique_ptr group_merge_m2(column_view const& values, using output_iterator = thrust::zip_iterator; auto const out_iter = output_iterator{cuda::std::make_tuple(result_counts->mutable_view().template data(), - result_means->mutable_view().template data(), - result_M2s->mutable_view().template data(), - validities.begin())}; + result_means->mutable_view().template data(), + result_M2s->mutable_view().template data(), + validities.begin())}; auto const count_valid = values.child(0); auto const mean_values = values.child(1); diff --git a/cpp/src/groupby/sort/group_replace_nulls.cu b/cpp/src/groupby/sort/group_replace_nulls.cu index 9d9611f5530..e755b4317ac 100644 --- a/cpp/src/groupby/sort/group_replace_nulls.cu +++ b/cpp/src/groupby/sort/group_replace_nulls.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -22,13 +22,13 @@ #include +#include #include #include #include #include #include #include -#include #include diff --git a/cpp/src/io/comp/nvcomp_adapter.cu b/cpp/src/io/comp/nvcomp_adapter.cu index 60c4d084d76..7cb83488191 100644 --- a/cpp/src/io/comp/nvcomp_adapter.cu +++ b/cpp/src/io/comp/nvcomp_adapter.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,9 +19,9 @@ #include +#include #include #include -#include namespace cudf::io::nvcomp { diff --git a/cpp/src/io/comp/statistics.cu b/cpp/src/io/comp/statistics.cu index 462622dcb1f..ac5dfee2e6f 100644 --- a/cpp/src/io/comp/statistics.cu +++ b/cpp/src/io/comp/statistics.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * Copyright (c) 2023-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/src/io/json/json_column.cu b/cpp/src/io/json/json_column.cu index 9a0b3a4f9e8..5223e329fe1 100644 --- a/cpp/src/io/json/json_column.cu +++ b/cpp/src/io/json/json_column.cu @@ -344,16 +344,17 @@ std::vector copy_strings_to_host(device_span input, rmm::device_uvector string_offsets(num_strings, stream); rmm::device_uvector string_lengths(num_strings, stream); auto d_offset_pairs = thrust::make_zip_iterator(node_range_begin.begin(), node_range_end.begin()); - thrust::transform(rmm::exec_policy(stream), - d_offset_pairs, - d_offset_pairs + num_strings, - thrust::make_zip_iterator(string_offsets.begin(), string_lengths.begin()), - [] __device__(auto const& offsets) { - // Note: first character for non-field columns - return cuda::std::make_tuple( - static_cast(cuda::std::get<0>(offsets)), - static_cast(cuda::std::get<1>(offsets) - cuda::std::get<0>(offsets))); - }); + thrust::transform( + rmm::exec_policy(stream), + d_offset_pairs, + d_offset_pairs + num_strings, + thrust::make_zip_iterator(string_offsets.begin(), string_lengths.begin()), + [] __device__(auto const& offsets) { + // Note: first character for non-field columns + return cuda::std::make_tuple( + static_cast(cuda::std::get<0>(offsets)), + static_cast(cuda::std::get<1>(offsets) - cuda::std::get<0>(offsets))); + }); cudf::io::parse_options_view options_view{}; options_view.quotechar = '\0'; // no quotes diff --git a/cpp/src/io/json/nested_json_gpu.cu b/cpp/src/io/json/nested_json_gpu.cu index 1ec2fe20094..1149828bc54 100644 --- a/cpp/src/io/json/nested_json_gpu.cu +++ b/cpp/src/io/json/nested_json_gpu.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -38,12 +38,12 @@ #include #include +#include #include #include #include #include #include -#include #include #include diff --git a/cpp/src/io/orc/stripe_enc.cu b/cpp/src/io/orc/stripe_enc.cu index b83fe9a006d..e5dc315639d 100644 --- a/cpp/src/io/orc/stripe_enc.cu +++ b/cpp/src/io/orc/stripe_enc.cu @@ -32,10 +32,10 @@ #include #include +#include #include #include #include -#include namespace cudf { namespace io { diff --git a/cpp/src/io/parquet/page_enc.cu b/cpp/src/io/parquet/page_enc.cu index 8637fdc6258..4bc7b3a4a8d 100644 --- a/cpp/src/io/parquet/page_enc.cu +++ b/cpp/src/io/parquet/page_enc.cu @@ -32,6 +32,7 @@ #include +#include #include #include #include @@ -41,7 +42,6 @@ #include #include #include -#include #include diff --git a/cpp/src/io/parquet/page_hdr.cu b/cpp/src/io/parquet/page_hdr.cu index 624c82639af..cc493ac0630 100644 --- a/cpp/src/io/parquet/page_hdr.cu +++ b/cpp/src/io/parquet/page_hdr.cu @@ -298,10 +298,11 @@ inline __device__ bool parse_header(cuda::std::tuple& op, byte_stre struct gpuParseDataPageHeader { __device__ bool operator()(byte_stream_s* bs) { - auto op = cuda::std::make_tuple(ParquetFieldInt32(1, bs->page.num_input_values), - ParquetFieldEnum(2, bs->page.encoding), - ParquetFieldEnum(3, bs->page.definition_level_encoding), - ParquetFieldEnum(4, bs->page.repetition_level_encoding)); + auto op = + cuda::std::make_tuple(ParquetFieldInt32(1, bs->page.num_input_values), + ParquetFieldEnum(2, bs->page.encoding), + ParquetFieldEnum(3, bs->page.definition_level_encoding), + ParquetFieldEnum(4, bs->page.repetition_level_encoding)); return parse_header(op, bs); } }; @@ -310,7 +311,7 @@ struct gpuParseDictionaryPageHeader { __device__ bool operator()(byte_stream_s* bs) { auto op = cuda::std::make_tuple(ParquetFieldInt32(1, bs->page.num_input_values), - ParquetFieldEnum(2, bs->page.encoding)); + ParquetFieldEnum(2, bs->page.encoding)); return parse_header(op, bs); } }; @@ -318,12 +319,13 @@ struct gpuParseDictionaryPageHeader { struct gpuParseDataPageHeaderV2 { __device__ bool operator()(byte_stream_s* bs) { - auto op = cuda::std::make_tuple(ParquetFieldInt32(1, bs->page.num_input_values), - ParquetFieldInt32(2, bs->page.num_nulls), - ParquetFieldInt32(3, bs->page.num_rows), - ParquetFieldEnum(4, bs->page.encoding), - ParquetFieldInt32(5, bs->page.lvl_bytes[level_type::DEFINITION]), - ParquetFieldInt32(6, bs->page.lvl_bytes[level_type::REPETITION])); + auto op = + cuda::std::make_tuple(ParquetFieldInt32(1, bs->page.num_input_values), + ParquetFieldInt32(2, bs->page.num_nulls), + ParquetFieldInt32(3, bs->page.num_rows), + ParquetFieldEnum(4, bs->page.encoding), + ParquetFieldInt32(5, bs->page.lvl_bytes[level_type::DEFINITION]), + ParquetFieldInt32(6, bs->page.lvl_bytes[level_type::REPETITION])); return parse_header(op, bs); } }; @@ -332,11 +334,11 @@ struct gpuParsePageHeader { __device__ bool operator()(byte_stream_s* bs) { auto op = cuda::std::make_tuple(ParquetFieldEnum(1, bs->page_type), - ParquetFieldInt32(2, bs->page.uncompressed_page_size), - ParquetFieldInt32(3, bs->page.compressed_page_size), - ParquetFieldStruct(5), - ParquetFieldStruct(7), - ParquetFieldStruct(8)); + ParquetFieldInt32(2, bs->page.uncompressed_page_size), + ParquetFieldInt32(3, bs->page.compressed_page_size), + ParquetFieldStruct(5), + ParquetFieldStruct(7), + ParquetFieldStruct(8)); return parse_header(op, bs); } }; diff --git a/cpp/src/io/parquet/reader_impl_helpers.hpp b/cpp/src/io/parquet/reader_impl_helpers.hpp index ea569fed494..dde28e11c11 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.hpp +++ b/cpp/src/io/parquet/reader_impl_helpers.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/src/io/text/bgzip_data_chunk_source.cu b/cpp/src/io/text/bgzip_data_chunk_source.cu index 199080609ac..94a48bcfb92 100644 --- a/cpp/src/io/text/bgzip_data_chunk_source.cu +++ b/cpp/src/io/text/bgzip_data_chunk_source.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -58,9 +58,9 @@ struct bgzip_nvcomp_transform_functor { auto const decompressed_begin = cuda::std::get<2>(t); auto const decompressed_end = cuda::std::get<3>(t); return cuda::std::make_tuple(device_span{compressed_ptr + compressed_begin, - compressed_end - compressed_begin}, - device_span{decompressed_ptr + decompressed_begin, - decompressed_end - decompressed_begin}); + compressed_end - compressed_begin}, + device_span{decompressed_ptr + decompressed_begin, + decompressed_end - decompressed_begin}); } }; diff --git a/cpp/src/io/utilities/string_parsing.hpp b/cpp/src/io/utilities/string_parsing.hpp index e3c95540888..9e8659768f1 100644 --- a/cpp/src/io/utilities/string_parsing.hpp +++ b/cpp/src/io/utilities/string_parsing.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * Copyright (c) 2023-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -22,8 +22,8 @@ #include -#include #include +#include namespace cudf::io { namespace detail { diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 3f647fd7a6a..57f4a8f636a 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -27,13 +27,13 @@ #include #include +#include #include #include #include #include #include #include -#include #include #include diff --git a/cpp/src/json/json_path.cu b/cpp/src/json/json_path.cu index a48423ed7eb..9d746d20b35 100644 --- a/cpp/src/json/json_path.cu +++ b/cpp/src/json/json_path.cu @@ -38,10 +38,10 @@ #include #include +#include #include #include #include -#include namespace cudf { namespace detail { diff --git a/cpp/src/lists/dremel.cu b/cpp/src/lists/dremel.cu index 51f708ba510..20fd70b7fb0 100644 --- a/cpp/src/lists/dremel.cu +++ b/cpp/src/lists/dremel.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -407,8 +407,8 @@ dremel_data get_encoding(column_view h_col, auto input_parent_zip_it = thrust::make_zip_iterator(cuda::std::make_tuple(input_parent_rep_it, input_parent_def_it)); - auto input_child_zip_it = - thrust::make_zip_iterator(cuda::std::make_tuple(temp_rep_vals.begin(), temp_def_vals.begin())); + auto input_child_zip_it = thrust::make_zip_iterator( + cuda::std::make_tuple(temp_rep_vals.begin(), temp_def_vals.begin())); auto output_zip_it = thrust::make_zip_iterator(cuda::std::make_tuple(rep_level.begin(), def_level.begin())); diff --git a/cpp/src/merge/merge.cu b/cpp/src/merge/merge.cu index 2d9bafc54c7..c07a00938b7 100644 --- a/cpp/src/merge/merge.cu +++ b/cpp/src/merge/merge.cu @@ -41,6 +41,7 @@ #include #include +#include #include #include #include @@ -48,7 +49,6 @@ #include #include #include -#include #include diff --git a/cpp/src/partitioning/round_robin.cu b/cpp/src/partitioning/round_robin.cu index c8b0f9b290f..aeaf5048ec5 100644 --- a/cpp/src/partitioning/round_robin.cu +++ b/cpp/src/partitioning/round_robin.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2023, NVIDIA CORPORATION. + * Copyright (c) 2019-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -31,6 +31,7 @@ #include #include +#include #include #include #include @@ -38,7 +39,6 @@ #include #include #include -#include #include diff --git a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu index 7450542af75..d39635751df 100644 --- a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu +++ b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu @@ -33,6 +33,7 @@ #include #include +#include #include #include #include @@ -50,7 +51,6 @@ #include #include #include -#include #include @@ -825,7 +825,7 @@ struct get_scalar_minmax { { return valid_count > 0 ? cuda::std::make_tuple(static_cast(col.element(0)), - static_cast(col.element(valid_count - 1))) + static_cast(col.element(valid_count - 1))) : cuda::std::make_tuple(0.0, 0.0); } }; @@ -867,7 +867,7 @@ struct typed_group_tdigest { thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + num_groups, thrust::make_zip_iterator(cuda::std::make_tuple(min_col->mutable_view().begin(), - max_col->mutable_view().begin())), + max_col->mutable_view().begin())), get_scalar_minmax_grouped{*d_col, group_offsets, group_valid_counts.begin()}); // for simple input values, the "centroids" all have a weight of 1. @@ -947,7 +947,7 @@ struct typed_reduce_tdigest { thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + 1, thrust::make_zip_iterator(cuda::std::make_tuple(min_col->mutable_view().begin(), - max_col->mutable_view().begin())), + max_col->mutable_view().begin())), get_scalar_minmax{*d_col, valid_count}); // for simple input values, the "centroids" all have a weight of 1. diff --git a/cpp/src/reductions/histogram.cu b/cpp/src/reductions/histogram.cu index 042706b9b4d..9ffafb72edd 100644 --- a/cpp/src/reductions/histogram.cu +++ b/cpp/src/reductions/histogram.cu @@ -21,9 +21,9 @@ #include #include +#include #include #include -#include #include #include diff --git a/cpp/src/replace/clamp.cu b/cpp/src/replace/clamp.cu index aabc1fc6656..60522fd7319 100644 --- a/cpp/src/replace/clamp.cu +++ b/cpp/src/replace/clamp.cu @@ -37,12 +37,12 @@ #include #include +#include #include #include #include #include #include -#include #include @@ -131,8 +131,8 @@ std::enable_if_t(), std::unique_ptr> clamp auto output_device_view = cudf::mutable_column_device_view::create(output->mutable_view(), stream); auto input_device_view = cudf::column_device_view::create(input, stream); - auto scalar_zip_itr = - thrust::make_zip_iterator(cuda::std::make_tuple(lo_itr, lo_replace_itr, hi_itr, hi_replace_itr)); + auto scalar_zip_itr = thrust::make_zip_iterator( + cuda::std::make_tuple(lo_itr, lo_replace_itr, hi_itr, hi_replace_itr)); auto trans = cuda::proclaim_return_type([] __device__(auto element_optional, auto scalar_tuple) { diff --git a/cpp/src/replace/nulls.cu b/cpp/src/replace/nulls.cu index 3a9aaf6aded..73fd3bd4060 100644 --- a/cpp/src/replace/nulls.cu +++ b/cpp/src/replace/nulls.cu @@ -43,6 +43,7 @@ #include #include +#include #include #include #include @@ -50,7 +51,6 @@ #include #include #include -#include namespace { // anonymous diff --git a/cpp/src/replace/replace.cu b/cpp/src/replace/replace.cu index 7b30926855c..2414db0865f 100644 --- a/cpp/src/replace/replace.cu +++ b/cpp/src/replace/replace.cu @@ -53,11 +53,11 @@ #include #include +#include #include #include #include #include -#include namespace { // anonymous diff --git a/cpp/src/sort/rank.cu b/cpp/src/sort/rank.cu index e50d10fd53f..3fbbfef5cb2 100644 --- a/cpp/src/sort/rank.cu +++ b/cpp/src/sort/rank.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -31,6 +31,7 @@ #include #include +#include #include #include #include @@ -42,7 +43,6 @@ #include #include #include -#include #include #include diff --git a/cpp/tests/io/type_inference_test.cu b/cpp/tests/io/type_inference_test.cu index 3b30ea1d965..a7282aba987 100644 --- a/cpp/tests/io/type_inference_test.cu +++ b/cpp/tests/io/type_inference_test.cu @@ -25,8 +25,8 @@ #include #include -#include #include +#include #include #include @@ -58,8 +58,8 @@ TEST_F(TypeInference, Basic) auto const d_string_length = cudf::detail::make_device_uvector_async( string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); - auto d_col_strings = - thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); + auto d_col_strings = thrust::make_zip_iterator( + cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -91,8 +91,8 @@ TEST_F(TypeInference, Null) auto const d_string_length = cudf::detail::make_device_uvector_async( string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); - auto d_col_strings = - thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); + auto d_col_strings = thrust::make_zip_iterator( + cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -124,8 +124,8 @@ TEST_F(TypeInference, AllNull) auto const d_string_length = cudf::detail::make_device_uvector_async( string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); - auto d_col_strings = - thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); + auto d_col_strings = thrust::make_zip_iterator( + cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -157,8 +157,8 @@ TEST_F(TypeInference, String) auto const d_string_length = cudf::detail::make_device_uvector_async( string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); - auto d_col_strings = - thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); + auto d_col_strings = thrust::make_zip_iterator( + cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -190,8 +190,8 @@ TEST_F(TypeInference, Bool) auto const d_string_length = cudf::detail::make_device_uvector_async( string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); - auto d_col_strings = - thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); + auto d_col_strings = thrust::make_zip_iterator( + cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -223,8 +223,8 @@ TEST_F(TypeInference, Timestamp) auto const d_string_length = cudf::detail::make_device_uvector_async( string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); - auto d_col_strings = - thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); + auto d_col_strings = thrust::make_zip_iterator( + cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), @@ -257,8 +257,8 @@ TEST_F(TypeInference, InvalidInput) auto const d_string_length = cudf::detail::make_device_uvector_async( string_length, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); - auto d_col_strings = - thrust::make_zip_iterator(cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); + auto d_col_strings = thrust::make_zip_iterator( + cuda::std::make_tuple(d_string_offset.begin(), d_string_length.begin())); auto res_type = infer_data_type(options.json_view(), diff --git a/cpp/tests/reductions/scan_tests.cpp b/cpp/tests/reductions/scan_tests.cpp index 45df895dbbf..f97faf3f57d 100644 --- a/cpp/tests/reductions/scan_tests.cpp +++ b/cpp/tests/reductions/scan_tests.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2023, NVIDIA CORPORATION. + * Copyright (c) 2019-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,9 +25,9 @@ #include #include +#include #include #include -#include #include #include diff --git a/cpp/tests/utilities/tdigest_utilities.cu b/cpp/tests/utilities/tdigest_utilities.cu index 193ec4ce7c5..de384d01c8d 100644 --- a/cpp/tests/utilities/tdigest_utilities.cu +++ b/cpp/tests/utilities/tdigest_utilities.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -26,10 +26,10 @@ #include +#include #include #include #include -#include // for use with groupby and reduction aggregation tests.