Remove floating point types from radix sort fast-path (#7215)

Closes #7212 

Reference #7167 (comment)

Using radix sort for all fixed-width types causes an [error in Spark when floating point columns contain NaN elements](NVIDIA/spark-rapids#1585).

This PR removes floating-point column types from the radix fast-path. This means the original `relational_compare` row operator is used to handle sorting floating point columns since they could possibly contain NaN elements.

The `NANSorting` gtest included null elements so it did not catch the fast-path output discrepancy. This PR adds a `NANSortingNonNull` gtest to check for the desired NaN sorting behavior.

Authors:
  - David (@davidwendt)

Approvers:
  - Jake Hemstad (@jrhemstad)
  - Conor Hoekstra (@codereport)

URL: #7215

Author: davidwendt

cpp/src/sort/sort_column.cu

@@ -24,6 +24,17 @@ namespace {
  * @brief Type-dispatched functor for sorting a single column.
  */
 struct column_sorted_order_fn {
+  /**
+   * @brief Compile time check for allowing radix sort for column type.
+   *
+   * Floating point is removed here for special handling of NaNs.
+   */
+  template <typename T>
+  static constexpr bool is_radix_sort_supported()
+  {
+    return cudf::is_fixed_width<T>() && !cudf::is_floating_point<T>();
+  }
+
   /**
    * @brief Sorts fixed-width columns using faster thrust sort.
    *
@@ -32,15 +43,15 @@ struct column_sorted_order_fn {
    * @param ascending True if sort order is ascending
    * @param stream CUDA stream used for device memory operations and kernel launches
    */
-  template <typename T, typename std::enable_if_t<cudf::is_fixed_width<T>()>* = nullptr>
+  template <typename T, typename std::enable_if_t<is_radix_sort_supported<T>()>* = nullptr>
   void radix_sort(column_view const& input,
                   mutable_column_view& indices,
                   bool ascending,
                   rmm::cuda_stream_view stream)
   {
-    // A non-stable sort on a fixed-width column with no nulls will use a radix sort
-    // if using only the thrust::less or thrust::greater comparators but also
-    // requires making a copy of the input data.
+    // A non-stable sort on a column of arithmetic type with no nulls will use a radix sort
+    // if specifying only the `thrust::less` or `thrust::greater` comparators.
+    // But this also requires making a copy of the input data.
     auto temp_col = column(input, stream);
     auto d_col    = temp_col.mutable_view();
     using DeviceT = device_storage_type_t<T>;
@@ -58,7 +69,7 @@ struct column_sorted_order_fn {
                           thrust::greater<DeviceT>());
     }
   }
-  template <typename T, typename std::enable_if_t<!cudf::is_fixed_width<T>()>* = nullptr>
+  template <typename T, typename std::enable_if_t<!is_radix_sort_supported<T>()>* = nullptr>
   void radix_sort(column_view const&, mutable_column_view&, bool, rmm::cuda_stream_view)
   {
     CUDF_FAIL("Only fixed-width types are suitable for faster sorting");
@@ -83,8 +94,8 @@ struct column_sorted_order_fn {
                   null_order null_precedence,
                   rmm::cuda_stream_view stream)
   {
-    // column with nulls or non-fixed-width column will also use a comparator
-    if (input.has_nulls() || !cudf::is_fixed_width<T>()) {
+    // column with nulls or non-supported types will also use a comparator
+    if (input.has_nulls() || !is_radix_sort_supported<T>()) {
       auto keys = column_device_view::create(input, stream);
       thrust::sort(rmm::exec_policy(stream),
                    indices.begin<size_type>(),

cpp/tests/table/row_operators_tests.cpp

@@ -65,3 +65,25 @@ TEST_F(RowOperatorTestForNAN, NANSorting)
 
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected2, got2->view());
 }
+
+TEST_F(RowOperatorTestForNAN, NANSortingNonNull)
+{
+  cudf::test::fixed_width_column_wrapper<double> input{
+    {0.,
+     double(NAN),
+     -1.,
+     7.,
+     std::numeric_limits<double>::infinity(),
+     1.,
+     -1 * std::numeric_limits<double>::infinity()}};
+
+  cudf::table_view input_table{{input}};
+
+  auto result = cudf::sorted_order(input_table, {cudf::order::ASCENDING});
+  cudf::test::fixed_width_column_wrapper<int32_t> expected_asc{{6, 2, 0, 5, 3, 4, 1}};
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_asc, result->view());
+
+  result = cudf::sorted_order(input_table, {cudf::order::DESCENDING});
+  cudf::test::fixed_width_column_wrapper<int32_t> expected_desc{{1, 4, 3, 5, 0, 2, 6}};
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_desc, result->view());
+}