From ba0febe308a8c097474b3316387dc8051fa1bc64 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Mon, 3 Oct 2022 17:29:42 -0700 Subject: [PATCH] Test/remove thrust vector usage (#11813) This PR removes usage of `thrust::device_vector` from almost all of our tests. Since the construction of a device vector is not stream-ordered, we should be using `rmm::device_uvector` instead wherever possible. There is one remaining use of `thrust::device_vector`, but that is in an test explicitly verifying that `device_vector` can convert implicitly to a `device_span` so it's worth keeping that there. I am working on automated tooling to detect any usage of stream 0 in tests as part of a push to prioritize stream-safety in libcudf, and this PR is a prerequisite to adding such tooling to our CI pipeline since at that point any test using stream 0 would fail. Since there is at least one test where I anticipate stream 0 will always be used (the one described above), I should be able to add specific tests to an allowlist as needed. It's an open question whether the added complexity required by the changes in this PR is a worthwhile tradeoff to be able to programmatically detect stream 0 usage. If reviewers feel that the additional complexity is too high, we can revert some (or all) of these changes and I can just plan for allowing stream 0 usage in all of the necessary tests. This PR demonstrates how we would go about removing it if we choose to do so, though. Authors: - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - Vukasin Milovanovic (https://github.com/vuule) - Tobias Ribizel (https://github.com/upsj) - Nghia Truong (https://github.com/ttnghia) URL: https://github.com/rapidsai/cudf/pull/11813 --- cpp/include/cudf_test/tdigest_utilities.cuh | 2 +- cpp/tests/join/conditional_join_tests.cu | 40 ++++++++++++++++----- cpp/tests/quantiles/tdigest_utilities.cu | 33 +++++++++++++---- cpp/tests/utilities_tests/span_tests.cu | 12 +++---- 4 files changed, 64 insertions(+), 23 deletions(-) diff --git a/cpp/include/cudf_test/tdigest_utilities.cuh b/cpp/include/cudf_test/tdigest_utilities.cuh index 6f206a789fd..1a75016d78c 100644 --- a/cpp/include/cudf_test/tdigest_utilities.cuh +++ b/cpp/include/cudf_test/tdigest_utilities.cuh @@ -118,7 +118,7 @@ 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( data_type{type_id::FLOAT64}, spans.size(), mask_state::UNALLOCATED); diff --git a/cpp/tests/join/conditional_join_tests.cu b/cpp/tests/join/conditional_join_tests.cu index bc2a96b5adf..f8dfc972191 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,30 @@ gen_random_nullable_repeated_columns(unsigned int N = 10000, unsigned int num_re std::pair(std::move(right), std::move(right_nulls))); } +// `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{}; + __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 +275,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 +286,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 +294,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( diff --git a/cpp/tests/quantiles/tdigest_utilities.cu b/cpp/tests/quantiles/tdigest_utilities.cu index 63ccd85bd6d..3cf2f2eb4ef 100644 --- a/cpp/tests/quantiles/tdigest_utilities.cu +++ b/cpp/tests/quantiles/tdigest_utilities.cu @@ -51,13 +51,34 @@ 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()); + 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 +86,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]; }); 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 {