From 7e0ae5156fa7d290542905e767f32b92a8c26530 Mon Sep 17 00:00:00 2001 From: David Wendt <45795991+davidwendt@users.noreply.github.com> Date: Tue, 20 Jun 2023 10:38:44 -0400 Subject: [PATCH] Fix memcheck errors found in REDUCTION_TEST (#13574) The nightly tests found memcheck errors in `REDUCTION_TEST` for **any/all** aggregation types on **dictionary** columns. The error is caused by the internal `atomicOr()` function which reads the target value as 32-bit even if the type is bool (8-bit). The fix works around this limitation by using `atomicOr()` on a 32-bit device scalar and then converting that to a bool scalar on return. All other tests in `REDUCTION_TEST` now pass memcheck successfully. Closes #13572 Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Nghia Truong (https://github.com/ttnghia) - https://github.com/nvdbaranec - Yunsong Wang (https://github.com/PointKernel) URL: https://github.com/rapidsai/cudf/pull/13574 --- cpp/src/reductions/all.cu | 8 +-- cpp/src/reductions/any.cu | 8 +-- cpp/tests/reductions/reduction_tests.cpp | 91 ++++++------------------ 3 files changed, 31 insertions(+), 76 deletions(-) diff --git a/cpp/src/reductions/all.cu b/cpp/src/reductions/all.cu index 9d32bc4c7f6..c2f278532ac 100644 --- a/cpp/src/reductions/all.cu +++ b/cpp/src/reductions/all.cu @@ -46,7 +46,7 @@ struct all_fn { if (*d_result && (iter[idx] != *d_result)) atomicAnd(d_result, false); } Iterator iter; - bool* d_result; + int32_t* d_result; }; template >* = nullptr> @@ -61,12 +61,12 @@ struct all_fn { cudf::dictionary::detail::make_dictionary_pair_iterator(*d_dict, input.has_nulls()); return thrust::make_transform_iterator(pair_iter, null_iter); }(); - auto result = std::make_unique>(true, true, stream, mr); + auto d_result = rmm::device_scalar(1, stream, rmm::mr::get_current_device_resource()); thrust::for_each_n(rmm::exec_policy(stream), thrust::make_counting_iterator(0), input.size(), - all_true_fn{iter, result->data()}); - return result; + all_true_fn{iter, d_result.data()}); + return std::make_unique>(d_result.value(stream), true, stream, mr); } template >* = nullptr> std::unique_ptr operator()(column_view const&, diff --git a/cpp/src/reductions/any.cu b/cpp/src/reductions/any.cu index 07977d2417f..d1b66d5c254 100644 --- a/cpp/src/reductions/any.cu +++ b/cpp/src/reductions/any.cu @@ -46,7 +46,7 @@ struct any_fn { if (!*d_result && (iter[idx] != *d_result)) atomicOr(d_result, true); } Iterator iter; - bool* d_result; + int32_t* d_result; }; template >* = nullptr> @@ -61,12 +61,12 @@ struct any_fn { cudf::dictionary::detail::make_dictionary_pair_iterator(*d_dict, input.has_nulls()); return thrust::make_transform_iterator(pair_iter, null_iter); }(); - auto result = std::make_unique>(false, true, stream, mr); + auto d_result = rmm::device_scalar(0, stream, rmm::mr::get_current_device_resource()); thrust::for_each_n(rmm::exec_policy(stream), thrust::make_counting_iterator(0), input.size(), - any_true_fn{iter, result->data()}); - return result; + any_true_fn{iter, d_result.data()}); + return std::make_unique>(d_result.value(stream), true, stream, mr); } template >* = nullptr> std::unique_ptr operator()(column_view const&, diff --git a/cpp/tests/reductions/reduction_tests.cpp b/cpp/tests/reductions/reduction_tests.cpp index dda0fd1f6f5..4c96f104ebb 100644 --- a/cpp/tests/reductions/reduction_tests.cpp +++ b/cpp/tests/reductions/reduction_tests.cpp @@ -2077,88 +2077,43 @@ TYPED_TEST(DictionaryAnyAllTest, AnyAll) std::vector v_some = convert_values(some_values); cudf::data_type output_dtype(cudf::type_id::BOOL8); + auto any_agg = cudf::make_any_aggregation(); + auto all_agg = cudf::make_all_aggregation(); + // without nulls { cudf::test::dictionary_column_wrapper all_col(v_all.begin(), v_all.end()); - EXPECT_TRUE(this - ->template reduction_test( - all_col, *cudf::make_any_aggregation(), output_dtype) - .first); - EXPECT_TRUE(this - ->template reduction_test( - all_col, *cudf::make_all_aggregation(), output_dtype) - .first); + EXPECT_TRUE(this->template reduction_test(all_col, *any_agg, output_dtype).first); + EXPECT_TRUE(this->template reduction_test(all_col, *all_agg, output_dtype).first); cudf::test::dictionary_column_wrapper none_col(v_none.begin(), v_none.end()); - EXPECT_FALSE(this - ->template reduction_test( - none_col, *cudf::make_any_aggregation(), output_dtype) - .first); - EXPECT_FALSE(this - ->template reduction_test( - none_col, *cudf::make_all_aggregation(), output_dtype) - .first); + EXPECT_FALSE(this->template reduction_test(none_col, *any_agg, output_dtype).first); + EXPECT_FALSE(this->template reduction_test(none_col, *all_agg, output_dtype).first); cudf::test::dictionary_column_wrapper some_col(v_some.begin(), v_some.end()); - EXPECT_TRUE(this - ->template reduction_test( - some_col, *cudf::make_any_aggregation(), output_dtype) - .first); - EXPECT_FALSE(this - ->template reduction_test( - some_col, *cudf::make_all_aggregation(), output_dtype) - .first); + EXPECT_TRUE(this->template reduction_test(some_col, *any_agg, output_dtype).first); + EXPECT_FALSE(this->template reduction_test(some_col, *all_agg, output_dtype).first); // sliced test - EXPECT_TRUE(this - ->template reduction_test(cudf::slice(some_col, {1, 3}).front(), - *cudf::make_any_aggregation(), - output_dtype) - .first); - EXPECT_TRUE(this - ->template reduction_test(cudf::slice(some_col, {1, 2}).front(), - *cudf::make_all_aggregation(), - output_dtype) - .first); + auto slice1 = cudf::slice(some_col, {1, 3}).front(); + auto slice2 = cudf::slice(some_col, {1, 2}).front(); + EXPECT_TRUE(this->template reduction_test(slice1, *any_agg, output_dtype).first); + EXPECT_TRUE(this->template reduction_test(slice2, *all_agg, output_dtype).first); } // with nulls { std::vector valid({1, 1, 0, 1}); cudf::test::dictionary_column_wrapper all_col(v_all.begin(), v_all.end(), valid.begin()); - EXPECT_TRUE(this - ->template reduction_test( - all_col, *cudf::make_any_aggregation(), output_dtype) - .first); - EXPECT_TRUE(this - ->template reduction_test( - all_col, *cudf::make_all_aggregation(), output_dtype) - .first); + EXPECT_TRUE(this->template reduction_test(all_col, *any_agg, output_dtype).first); + EXPECT_TRUE(this->template reduction_test(all_col, *all_agg, output_dtype).first); cudf::test::dictionary_column_wrapper none_col(v_none.begin(), v_none.end(), valid.begin()); - EXPECT_FALSE(this - ->template reduction_test( - none_col, *cudf::make_any_aggregation(), output_dtype) - .first); - EXPECT_FALSE(this - ->template reduction_test( - none_col, *cudf::make_all_aggregation(), output_dtype) - .first); + EXPECT_FALSE(this->template reduction_test(none_col, *any_agg, output_dtype).first); + EXPECT_FALSE(this->template reduction_test(none_col, *all_agg, output_dtype).first); cudf::test::dictionary_column_wrapper some_col(v_some.begin(), v_some.end(), valid.begin()); - EXPECT_TRUE(this - ->template reduction_test( - some_col, *cudf::make_any_aggregation(), output_dtype) - .first); - EXPECT_FALSE(this - ->template reduction_test( - some_col, *cudf::make_all_aggregation(), output_dtype) - .first); + EXPECT_TRUE(this->template reduction_test(some_col, *any_agg, output_dtype).first); + EXPECT_FALSE(this->template reduction_test(some_col, *all_agg, output_dtype).first); // sliced test - EXPECT_TRUE(this - ->template reduction_test(cudf::slice(some_col, {0, 3}).front(), - *cudf::make_any_aggregation(), - output_dtype) - .first); - EXPECT_TRUE(this - ->template reduction_test(cudf::slice(some_col, {1, 4}).front(), - *cudf::make_all_aggregation(), - output_dtype) - .first); + auto slice1 = cudf::slice(some_col, {0, 3}).front(); + auto slice2 = cudf::slice(some_col, {1, 4}).front(); + EXPECT_TRUE(this->template reduction_test(slice1, *any_agg, output_dtype).first); + EXPECT_TRUE(this->template reduction_test(slice2, *all_agg, output_dtype).first); } }