From 9ade81e47955fea401a03e60c5681d1f1d524f1d Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 22 Sep 2022 16:37:50 -0700 Subject: [PATCH 1/6] Fix stream usage in tdigest tests (unclear if this is the right path forward though since it requires replacing all uses of device_vector with device_uvector in tests, which implies that users will have to do the same). --- cpp/include/cudf_test/tdigest_utilities.cuh | 73 +++++++++++---------- cpp/tests/quantiles/tdigest_utilities.cu | 37 +++++++++-- 2 files changed, 68 insertions(+), 42 deletions(-) diff --git a/cpp/include/cudf_test/tdigest_utilities.cuh b/cpp/include/cudf_test/tdigest_utilities.cuh index 6f206a789fd..f23d7ccf58f 100644 --- a/cpp/include/cudf_test/tdigest_utilities.cuh +++ b/cpp/include/cudf_test/tdigest_utilities.cuh @@ -118,7 +118,8 @@ void tdigest_minmax_compare(cudf::tdigest::tdigest_column_view const& tdv, // verify min/max thrust::host_vector> h_spans; h_spans.push_back({input_values.begin(), static_cast(input_values.size())}); - thrust::device_vector> spans(h_spans); + // thrust::device_vector> spans(h_spans); + auto spans = cudf::detail::make_device_uvector_async(h_spans, cudf::default_stream_value); auto expected_min = cudf::make_fixed_width_column( data_type{type_id::FLOAT64}, spans.size(), mask_state::UNALLOCATED); @@ -259,41 +260,41 @@ void tdigest_simple_large_input_double_aggregation(Func op) } // delta 100 - { - int const delta = 100; - auto result = cudf::type_dispatcher(values->view().type(), tdigest_gen{}, op, *values, delta); - std::vector expected{{0, 0.07265722021410986331, 739}, - {7, 8.19766194442652640362, 10693}, - {16, 36.82277869518204482802, 20276}, - {29, 72.95424834129075009059, 22623}, - {38, 90.61229683516096145013, 15581}, - {46, 99.07283498858802772702, 5142}, - {50, 99.99970905482754801596, 1}}; - cudf::tdigest::tdigest_column_view tdv(*result); - - tdigest_sample_compare(tdv, expected); - - // verify min/max - tdigest_minmax_compare(tdv, *values); - } - - // delta 10 - { - int const delta = 10; - auto result = cudf::type_dispatcher(values->view().type(), tdigest_gen{}, op, *values, delta); - std::vector expected{{0, 7.15508346777729631327, 71618}, - {1, 33.04971680740474226923, 187499}, - {2, 62.50566666553867634093, 231762}, - {3, 83.46216572053654658703, 187500}, - {4, 96.42204425201593664951, 71620}, - {5, 99.99970905482754801596, 1}}; - cudf::tdigest::tdigest_column_view tdv(*result); - - tdigest_sample_compare(tdv, expected); - - // verify min/max - tdigest_minmax_compare(tdv, *values); - } + //{ + // int const delta = 100; + // auto result = cudf::type_dispatcher(values->view().type(), tdigest_gen{}, op, *values, delta); + // std::vector expected{{0, 0.07265722021410986331, 739}, + // {7, 8.19766194442652640362, 10693}, + // {16, 36.82277869518204482802, 20276}, + // {29, 72.95424834129075009059, 22623}, + // {38, 90.61229683516096145013, 15581}, + // {46, 99.07283498858802772702, 5142}, + // {50, 99.99970905482754801596, 1}}; + // cudf::tdigest::tdigest_column_view tdv(*result); + // + // tdigest_sample_compare(tdv, expected); + // + // // verify min/max + // tdigest_minmax_compare(tdv, *values); + //} + // + //// delta 10 + //{ + // int const delta = 10; + // auto result = cudf::type_dispatcher(values->view().type(), tdigest_gen{}, op, *values, delta); + // std::vector expected{{0, 7.15508346777729631327, 71618}, + // {1, 33.04971680740474226923, 187499}, + // {2, 62.50566666553867634093, 231762}, + // {3, 83.46216572053654658703, 187500}, + // {4, 96.42204425201593664951, 71620}, + // {5, 99.99970905482754801596, 1}}; + // cudf::tdigest::tdigest_column_view tdv(*result); + // + // tdigest_sample_compare(tdv, expected); + // + // // verify min/max + // tdigest_minmax_compare(tdv, *values); + //} } // shared test for groupby/reduction. diff --git a/cpp/tests/quantiles/tdigest_utilities.cu b/cpp/tests/quantiles/tdigest_utilities.cu index 63ccd85bd6d..922bf62d70a 100644 --- a/cpp/tests/quantiles/tdigest_utilities.cu +++ b/cpp/tests/quantiles/tdigest_utilities.cu @@ -51,13 +51,38 @@ void tdigest_sample_compare(cudf::tdigest::tdigest_column_view const& tdv, auto sampled_result_weight = cudf::make_fixed_width_column( data_type{type_id::FLOAT64}, h_expected.size(), mask_state::UNALLOCATED); - rmm::device_vector expected(h_expected.begin(), h_expected.end()); + // We have to unpack the vector> into three separate vectors for + // copying into device uvectors because rmm::device_uvector doesn't support + // thrust::tuple (it's not trivially copyable). + // TODO: thrust::tuple_element should work here, but doesn't seem to play nice + auto h_expected_src = std::vector(h_expected.size()); + auto h_expected_mean = std::vector(h_expected.size()); + auto h_expected_weight = std::vector(h_expected.size()); + + { + 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]); + }); + } + + auto d_expected_src = + cudf::detail::make_device_uvector_async(h_expected_src, cudf::default_stream_value); + auto d_expected_mean = + cudf::detail::make_device_uvector_async(h_expected_mean, cudf::default_stream_value); + auto d_expected_weight = + cudf::detail::make_device_uvector_async(h_expected_weight, cudf::default_stream_value); + auto iter = thrust::make_counting_iterator(0); thrust::for_each( rmm::exec_policy(cudf::default_stream_value), iter, - iter + expected.size(), - [expected = expected.data().get(), + iter + h_expected.size(), + [expected_src_in = d_expected_src.data(), + expected_mean_in = d_expected_mean.data(), + expected_weight_in = d_expected_weight.data(), expected_mean = expected_mean->mutable_view().begin(), expected_weight = expected_weight->mutable_view().begin(), result_mean = result_mean.begin(), @@ -65,9 +90,9 @@ void tdigest_sample_compare(cudf::tdigest::tdigest_column_view const& tdv, sampled_result_mean = sampled_result_mean->mutable_view().begin(), sampled_result_weight = sampled_result_weight->mutable_view().begin()] __device__(size_type index) { - expected_mean[index] = thrust::get<1>(expected[index]); - expected_weight[index] = thrust::get<2>(expected[index]); - auto const src_index = thrust::get<0>(expected[index]); + expected_mean[index] = expected_mean_in[index]; + expected_weight[index] = expected_weight_in[index]; + auto const src_index = expected_src_in[index]; sampled_result_mean[index] = result_mean[src_index]; sampled_result_weight[index] = result_weight[src_index]; }); From b6445bcb3c005fb005497d4cb0f0677e40333002 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 28 Sep 2022 12:31:29 -0700 Subject: [PATCH 2/6] Remove device_vector usage in conditional join tests to remove default stream usage. --- cpp/tests/join/conditional_join_tests.cu | 36 ++++++++++++++++++------ 1 file changed, 27 insertions(+), 9 deletions(-) diff --git a/cpp/tests/join/conditional_join_tests.cu b/cpp/tests/join/conditional_join_tests.cu index bc2a96b5adf..872ffff8f08 100644 --- a/cpp/tests/join/conditional_join_tests.cu +++ b/cpp/tests/join/conditional_join_tests.cu @@ -26,10 +26,8 @@ #include -#include #include -#include -#include +#include #include #include @@ -127,6 +125,26 @@ gen_random_nullable_repeated_columns(unsigned int N = 10000, unsigned int num_re std::pair(std::move(right), std::move(right_nulls))); } +// Basic trivially copyable replacement for thrust::pair for use with `device_uvector` +struct index_pair { + cudf::size_type first{}; + cudf::size_type second{}; + __device__ index_pair(){}; + __device__ index_pair(cudf::size_type const& first, cudf::size_type const& second) + : first(first), second(second){}; +}; + +__device__ inline bool operator<(const index_pair& lhs, const index_pair& rhs) +{ + if (lhs.first > rhs.first) return false; + return (lhs.first < rhs.first) || (lhs.second < rhs.second); +} + +__device__ inline bool operator==(const index_pair& lhs, const index_pair& rhs) +{ + return lhs.first == rhs.first && lhs.second == rhs.second; +} + } // namespace /** @@ -253,10 +271,10 @@ struct ConditionalJoinPairReturnTest : public ConditionalJoinTest { */ void _compare_to_hash_join(PairJoinReturn const& result, PairJoinReturn const& reference) { - thrust::device_vector> result_pairs( - result.first->size()); - thrust::device_vector> reference_pairs( - reference.first->size()); + auto result_pairs = + rmm::device_uvector(result.first->size(), cudf::default_stream_value); + auto reference_pairs = + rmm::device_uvector(reference.first->size(), cudf::default_stream_value); thrust::transform(rmm::exec_policy(cudf::default_stream_value), result.first->begin(), @@ -264,7 +282,7 @@ struct ConditionalJoinPairReturnTest : public ConditionalJoinTest { result.second->begin(), result_pairs.begin(), [] __device__(cudf::size_type first, cudf::size_type second) { - return thrust::make_pair(first, second); + return index_pair{first, second}; }); thrust::transform(rmm::exec_policy(cudf::default_stream_value), reference.first->begin(), @@ -272,7 +290,7 @@ struct ConditionalJoinPairReturnTest : public ConditionalJoinTest { reference.second->begin(), reference_pairs.begin(), [] __device__(cudf::size_type first, cudf::size_type second) { - return thrust::make_pair(first, second); + return index_pair{first, second}; }); thrust::sort( From fe7c0051ef0e7a84bde8f6a215a4c83b4b417d28 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 28 Sep 2022 16:13:49 -0700 Subject: [PATCH 3/6] Remove device_vector usage from one span test. --- cpp/tests/utilities_tests/span_tests.cu | 12 +++++------- 1 file changed, 5 insertions(+), 7 deletions(-) diff --git a/cpp/tests/utilities_tests/span_tests.cu b/cpp/tests/utilities_tests/span_tests.cu index fc4104c765b..cccef4b6284 100644 --- a/cpp/tests/utilities_tests/span_tests.cu +++ b/cpp/tests/utilities_tests/span_tests.cu @@ -14,6 +14,7 @@ * limitations under the License. */ +#include #include #include #include @@ -234,17 +235,14 @@ __global__ void simple_device_kernel(device_span result) { result[0] = tru TEST(SpanTest, CanUseDeviceSpan) { - rmm::device_vector d_message = std::vector({false}); + auto d_message = + cudf::detail::make_zeroed_device_uvector_async(1, cudf::default_stream_value); - auto d_span = device_span(d_message.data().get(), d_message.size()); + auto d_span = device_span(d_message.data(), d_message.size()); simple_device_kernel<<<1, 1, 0, cudf::default_stream_value.value()>>>(d_span); - cudaDeviceSynchronize(); - - thrust::host_vector h_message = d_message; - - ASSERT_TRUE(h_message[0]); + ASSERT_TRUE(d_message.element(0, cudf::default_stream_value)); } class MdSpanTest : public cudf::test::BaseFixture { From 22735f4e7c7f5a423182c4673a8527796cfc14ff Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 28 Sep 2022 16:16:11 -0700 Subject: [PATCH 4/6] Remove outdated comments. --- cpp/include/cudf_test/tdigest_utilities.cuh | 71 ++++++++++----------- 1 file changed, 35 insertions(+), 36 deletions(-) diff --git a/cpp/include/cudf_test/tdigest_utilities.cuh b/cpp/include/cudf_test/tdigest_utilities.cuh index f23d7ccf58f..1a75016d78c 100644 --- a/cpp/include/cudf_test/tdigest_utilities.cuh +++ b/cpp/include/cudf_test/tdigest_utilities.cuh @@ -118,7 +118,6 @@ void tdigest_minmax_compare(cudf::tdigest::tdigest_column_view const& tdv, // verify min/max thrust::host_vector> h_spans; h_spans.push_back({input_values.begin(), static_cast(input_values.size())}); - // thrust::device_vector> spans(h_spans); auto spans = cudf::detail::make_device_uvector_async(h_spans, cudf::default_stream_value); auto expected_min = cudf::make_fixed_width_column( @@ -260,41 +259,41 @@ void tdigest_simple_large_input_double_aggregation(Func op) } // delta 100 - //{ - // int const delta = 100; - // auto result = cudf::type_dispatcher(values->view().type(), tdigest_gen{}, op, *values, delta); - // std::vector expected{{0, 0.07265722021410986331, 739}, - // {7, 8.19766194442652640362, 10693}, - // {16, 36.82277869518204482802, 20276}, - // {29, 72.95424834129075009059, 22623}, - // {38, 90.61229683516096145013, 15581}, - // {46, 99.07283498858802772702, 5142}, - // {50, 99.99970905482754801596, 1}}; - // cudf::tdigest::tdigest_column_view tdv(*result); - // - // tdigest_sample_compare(tdv, expected); - // - // // verify min/max - // tdigest_minmax_compare(tdv, *values); - //} - // - //// delta 10 - //{ - // int const delta = 10; - // auto result = cudf::type_dispatcher(values->view().type(), tdigest_gen{}, op, *values, delta); - // std::vector expected{{0, 7.15508346777729631327, 71618}, - // {1, 33.04971680740474226923, 187499}, - // {2, 62.50566666553867634093, 231762}, - // {3, 83.46216572053654658703, 187500}, - // {4, 96.42204425201593664951, 71620}, - // {5, 99.99970905482754801596, 1}}; - // cudf::tdigest::tdigest_column_view tdv(*result); - // - // tdigest_sample_compare(tdv, expected); - // - // // verify min/max - // tdigest_minmax_compare(tdv, *values); - //} + { + int const delta = 100; + auto result = cudf::type_dispatcher(values->view().type(), tdigest_gen{}, op, *values, delta); + std::vector expected{{0, 0.07265722021410986331, 739}, + {7, 8.19766194442652640362, 10693}, + {16, 36.82277869518204482802, 20276}, + {29, 72.95424834129075009059, 22623}, + {38, 90.61229683516096145013, 15581}, + {46, 99.07283498858802772702, 5142}, + {50, 99.99970905482754801596, 1}}; + cudf::tdigest::tdigest_column_view tdv(*result); + + tdigest_sample_compare(tdv, expected); + + // verify min/max + tdigest_minmax_compare(tdv, *values); + } + + // delta 10 + { + int const delta = 10; + auto result = cudf::type_dispatcher(values->view().type(), tdigest_gen{}, op, *values, delta); + std::vector expected{{0, 7.15508346777729631327, 71618}, + {1, 33.04971680740474226923, 187499}, + {2, 62.50566666553867634093, 231762}, + {3, 83.46216572053654658703, 187500}, + {4, 96.42204425201593664951, 71620}, + {5, 99.99970905482754801596, 1}}; + cudf::tdigest::tdigest_column_view tdv(*result); + + tdigest_sample_compare(tdv, expected); + + // verify min/max + tdigest_minmax_compare(tdv, *values); + } } // shared test for groupby/reduction. From c389c8f2e9072188a2bf69378f243cb4868f4eaf Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 28 Sep 2022 16:55:42 -0700 Subject: [PATCH 5/6] Remove outdated comments. --- cpp/tests/quantiles/tdigest_utilities.cu | 4 ---- 1 file changed, 4 deletions(-) diff --git a/cpp/tests/quantiles/tdigest_utilities.cu b/cpp/tests/quantiles/tdigest_utilities.cu index 922bf62d70a..3cf2f2eb4ef 100644 --- a/cpp/tests/quantiles/tdigest_utilities.cu +++ b/cpp/tests/quantiles/tdigest_utilities.cu @@ -51,10 +51,6 @@ void tdigest_sample_compare(cudf::tdigest::tdigest_column_view const& tdv, auto sampled_result_weight = cudf::make_fixed_width_column( data_type{type_id::FLOAT64}, h_expected.size(), mask_state::UNALLOCATED); - // We have to unpack the vector> into three separate vectors for - // copying into device uvectors because rmm::device_uvector doesn't support - // thrust::tuple (it's not trivially copyable). - // TODO: thrust::tuple_element should work here, but doesn't seem to play nice auto h_expected_src = std::vector(h_expected.size()); auto h_expected_mean = std::vector(h_expected.size()); auto h_expected_weight = std::vector(h_expected.size()); From b4724dc49e82fc7781d12878df4e4b76a50ffa19 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Mon, 3 Oct 2022 15:13:47 -0700 Subject: [PATCH 6/6] Add longer description. --- cpp/tests/join/conditional_join_tests.cu | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/cpp/tests/join/conditional_join_tests.cu b/cpp/tests/join/conditional_join_tests.cu index 872ffff8f08..f8dfc972191 100644 --- a/cpp/tests/join/conditional_join_tests.cu +++ b/cpp/tests/join/conditional_join_tests.cu @@ -125,7 +125,11 @@ gen_random_nullable_repeated_columns(unsigned int N = 10000, unsigned int num_re std::pair(std::move(right), std::move(right_nulls))); } -// Basic trivially copyable replacement for thrust::pair for use with `device_uvector` +// `rmm::device_uvector` requires that T be trivially copyable. `thrust::pair` does +// not satisfy this requirement because it defines nontrivial copy/move +// constructors. Therefore, we need a simple, trivially copyable pair-like +// object. `index_pair` is a minimal implementation suitable for use in the +// tests in this file. struct index_pair { cudf::size_type first{}; cudf::size_type second{};