From 60a88fb1a6ac2416bf0a922d70d4bec17c1491f7 Mon Sep 17 00:00:00 2001 From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com> Date: Fri, 26 Nov 2021 01:55:40 +0530 Subject: [PATCH 1/5] add __host__ for nvcc return type deduction Fixes #9703 add __host__ for nvcc return type deduction to work. replaced `auto` (generic lambda) with size_type. --- cpp/src/sort/rank.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/sort/rank.cu b/cpp/src/sort/rank.cu index c8a908e44cd..d961cd2cbad 100644 --- a/cpp/src/sort/rank.cu +++ b/cpp/src/sort/rank.cu @@ -209,7 +209,7 @@ void rank_average(cudf::device_span group_keys, tie_break_ranks_transform( group_keys, cudf::detail::make_counting_transform_iterator(1, - [] __device__(auto i) { + [] __host__ __device__(size_type i) -> MinCount { return MinCount{i, 1}; }), sorted_order_view, From d88fb12040f51405df087d4030f652e803fdb40d Mon Sep 17 00:00:00 2001 From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com> Date: Tue, 30 Nov 2021 09:44:58 +0530 Subject: [PATCH 2/5] replace host device lambda with device functor replace host device lambda with device functor due to device lambda return type deduction limitation --- cpp/src/sort/rank.cu | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/cpp/src/sort/rank.cu b/cpp/src/sort/rank.cu index d961cd2cbad..9056f22b3f0 100644 --- a/cpp/src/sort/rank.cu +++ b/cpp/src/sort/rank.cu @@ -194,6 +194,12 @@ void rank_max(cudf::device_span group_keys, stream); } +// Returns index, count +template +struct index_counter { + __device__ T operator()(size_type i) { return T{i, 1}; } +}; + void rank_average(cudf::device_span group_keys, column_view sorted_order_view, mutable_column_view rank_mutable_view, @@ -208,10 +214,7 @@ void rank_average(cudf::device_span group_keys, using MinCount = thrust::pair; tie_break_ranks_transform( group_keys, - cudf::detail::make_counting_transform_iterator(1, - [] __host__ __device__(size_type i) -> MinCount { - return MinCount{i, 1}; - }), + cudf::detail::make_counting_transform_iterator(1, index_counter{}), sorted_order_view, rank_mutable_view.begin(), [] __device__(auto rank_count1, auto rank_count2) { From 2ff54e3aa79b1237777a6e51c72dc7ab10c866b8 Mon Sep 17 00:00:00 2001 From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com> Date: Tue, 30 Nov 2021 23:34:47 +0530 Subject: [PATCH 3/5] [skip-ci] add comment --- cpp/src/sort/rank.cu | 2 ++ 1 file changed, 2 insertions(+) diff --git a/cpp/src/sort/rank.cu b/cpp/src/sort/rank.cu index 9056f22b3f0..c95aefc350b 100644 --- a/cpp/src/sort/rank.cu +++ b/cpp/src/sort/rank.cu @@ -214,6 +214,8 @@ void rank_average(cudf::device_span group_keys, using MinCount = thrust::pair; tie_break_ranks_transform( group_keys, + // use device functor with return type. Cannot use device lambda due to limitation. + // https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#extended-lambda-restrictions cudf::detail::make_counting_transform_iterator(1, index_counter{}), sorted_order_view, rank_mutable_view.begin(), From 0682a1f75c1a2645dfffe6c7317a65b94bab28a3 Mon Sep 17 00:00:00 2001 From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com> Date: Wed, 1 Dec 2021 12:27:05 +0530 Subject: [PATCH 4/5] add unit test for crash scenario --- cpp/tests/sort/rank_test.cpp | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/cpp/tests/sort/rank_test.cpp b/cpp/tests/sort/rank_test.cpp index 94e389fc7ce..926ad1e203e 100644 --- a/cpp/tests/sort/rank_test.cpp +++ b/cpp/tests/sort/rank_test.cpp @@ -410,5 +410,19 @@ TYPED_TEST(Rank, min_desc_bottom_pct) this->run_all_tests(rank_method::MIN, desc_bottom, col1_rank, col2_rank, col3_rank, true); } +struct RankLarge : public BaseFixture { +}; + +TEST_F(RankLarge, average_large) +{ + // testcase of https://github.com/rapidsai/cudf/issues/9703 + auto iter = thrust::counting_iterator(0); + fixed_width_column_wrapper col1(iter, iter + 10558); + auto result = + cudf::rank(col1, rank_method::AVERAGE, {}, null_policy::EXCLUDE, null_order::AFTER, false); + fixed_width_column_wrapper expected(iter + 1, iter + 10559); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(result->view(), expected); +} + } // namespace test } // namespace cudf From eb166db549a9cf89a497ec4d98741f3489de58b8 Mon Sep 17 00:00:00 2001 From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com> Date: Wed, 1 Dec 2021 20:28:16 +0530 Subject: [PATCH 5/5] Update cpp/src/sort/rank.cu Co-authored-by: David Wendt <45795991+davidwendt@users.noreply.github.com> --- cpp/src/sort/rank.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/sort/rank.cu b/cpp/src/sort/rank.cu index c95aefc350b..73a6ba6030c 100644 --- a/cpp/src/sort/rank.cu +++ b/cpp/src/sort/rank.cu @@ -214,7 +214,7 @@ void rank_average(cudf::device_span group_keys, using MinCount = thrust::pair; tie_break_ranks_transform( group_keys, - // use device functor with return type. Cannot use device lambda due to limitation. + // Use device functor with return type. Cannot use device lambda due to limitation. // https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#extended-lambda-restrictions cudf::detail::make_counting_transform_iterator(1, index_counter{}), sorted_order_view,