[EXPER] Add single column strings sort fast-path

cpp/CMakeLists.txt

@@ -1080,6 +1080,7 @@ add_library(
   src/sort/sort_column.cu
   src/sort/sort_radix.cu
   src/sort/sorted_order_radix.cu
+  src/sort/sorted_order_strings.cu
   src/sort/stable_sort_column.cu
   src/sort/stable_sort.cu
   src/sort/top_k.cu

cpp/benchmarks/sort/sort_strings.cpp

@@ -22,7 +22,7 @@ static void bench_sort_strings(nvbench::state& state)
   data_profile const profile = data_profile_builder().distribution(
     cudf::type_id::STRING, distribution_id::NORMAL, min_width, max_width);
 
-  auto const table = create_random_table({cudf::type_id::STRING}, row_count{num_rows});
+  auto const table = create_random_table({cudf::type_id::STRING}, row_count{num_rows}, profile);
 
   auto sv    = cudf::strings_column_view(table->view().column(0));
   auto bytes = sv.chars_size(cudf::get_default_stream());
@@ -44,3 +44,75 @@ NVBENCH_BENCH(bench_sort_strings)
   .add_int64_axis("min_width", {0})
   .add_int64_axis("max_width", {32, 64, 128, 256})
   .add_int64_axis("num_rows", {32768, 262144, 2097152});
+
+// Measures the `sorted_order` fast-path case: a single strings column with no nulls
+static void bench_sorted_order_strings(nvbench::state& state)
+{
+  auto const num_rows  = static_cast<cudf::size_type>(state.get_int64("num_rows"));
+  auto const min_width = static_cast<cudf::size_type>(state.get_int64("min_width"));
+  auto const max_width = static_cast<cudf::size_type>(state.get_int64("max_width"));
+
+  data_profile const profile =
+    data_profile_builder()
+      .distribution(cudf::type_id::STRING, distribution_id::NORMAL, min_width, max_width)
+      .no_validity();
+
+  auto const table = create_random_table({cudf::type_id::STRING}, row_count{num_rows}, profile);
+
+  auto sv    = cudf::strings_column_view(table->view().column(0));
+  auto bytes = sv.chars_size(cudf::get_default_stream());
+
+  state.set_cuda_stream(nvbench::make_cuda_stream_view(cudf::get_default_stream().get()));
+  state.add_global_memory_reads<nvbench::int8_t>(bytes);
+  state.add_global_memory_writes<cudf::size_type>(num_rows);
+
+  auto const mem_stats_logger = cudf::memory_stats_logger();
+
+  state.exec(nvbench::exec_tag::sync,
+             [&](nvbench::launch& launch) { cudf::sorted_order(table->view()); });
+
+  state.add_buffer_size(
+    mem_stats_logger.peak_memory_usage(), "peak_memory_usage", "peak_memory_usage");
+}
+
+NVBENCH_BENCH(bench_sorted_order_strings)
+  .set_name("sorted_order_strings")
+  .add_int64_axis("min_width", {1})
+  .add_int64_axis("max_width", {8, 32, 64, 128, 256})
+  .add_int64_axis("num_rows", {32768, 262144, 2097152, 16777216});
+
+// Measures the `stable_sorted_order` fast-path case: a single strings column with no nulls
+static void bench_stable_sorted_order_strings(nvbench::state& state)
+{
+  auto const num_rows  = static_cast<cudf::size_type>(state.get_int64("num_rows"));
+  auto const min_width = static_cast<cudf::size_type>(state.get_int64("min_width"));
+  auto const max_width = static_cast<cudf::size_type>(state.get_int64("max_width"));
+
+  data_profile const profile =
+    data_profile_builder()
+      .distribution(cudf::type_id::STRING, distribution_id::NORMAL, min_width, max_width)
+      .no_validity();
+
+  auto const table = create_random_table({cudf::type_id::STRING}, row_count{num_rows}, profile);
+
+  auto sv    = cudf::strings_column_view(table->view().column(0));
+  auto bytes = sv.chars_size(cudf::get_default_stream());
+
+  state.set_cuda_stream(nvbench::make_cuda_stream_view(cudf::get_default_stream().get()));
+  state.add_global_memory_reads<nvbench::int8_t>(bytes);
+  state.add_global_memory_writes<cudf::size_type>(num_rows);
+
+  auto const mem_stats_logger = cudf::memory_stats_logger();
+
+  state.exec(nvbench::exec_tag::sync,
+             [&](nvbench::launch& launch) { cudf::stable_sorted_order(table->view()); });
+
+  state.add_buffer_size(
+    mem_stats_logger.peak_memory_usage(), "peak_memory_usage", "peak_memory_usage");
+}
+
+NVBENCH_BENCH(bench_stable_sorted_order_strings)
+  .set_name("stable_sorted_order_strings")
+  .add_int64_axis("min_width", {1})
+  .add_int64_axis("max_width", {8, 32, 64, 128, 256})
+  .add_int64_axis("num_rows", {32768, 262144, 2097152, 16777216});

cpp/src/sort/sort_column.cu

@@ -5,6 +5,7 @@
 
 #include "sort_column_impl.cuh"
 #include "sort_radix.hpp"
+#include "sort_strings.hpp"
 
 #include <cudf/column/column_factories.hpp>
 #include <cudf/column/column_view.hpp>
@@ -31,6 +32,9 @@ std::unique_ptr<column> sorted_order<sort_method::UNSTABLE>(column_view const& i
   mutable_column_view indices_view = sorted_indices->mutable_view();
   if (is_radix_sortable(input)) {
     sorted_order_radix(input, indices_view, column_order == order::ASCENDING, stream);
+  } else if (is_strings_sortable(input)) {
+    sorted_order_strings<sort_method::UNSTABLE>(
+      input, indices_view, column_order == order::ASCENDING, stream);
   } else {
     cudf::type_dispatcher<dispatch_storage_type>(input.type(),
                                                  column_sorted_order_fn<sort_method::UNSTABLE>{},

cpp/src/sort/sort_strings.hpp

@@ -0,0 +1,43 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include "sort.hpp"
+
+#include <cudf/column/column_view.hpp>
+#include <cudf/types.hpp>
+
+#include <cuda/stream>
+
+namespace cudf {
+namespace detail {
+
+/**
+ * @brief Check if the strings fast-path sort is available for the given column
+ *
+ * @param column The column to check
+ * @return true if the strings fast-path sort is available, false otherwise
+ */
+bool is_strings_sortable(column_view const& column);
+
+/**
+ * @brief Sort indices of a single strings column
+ *
+ * This should only be used for non-empty strings columns with no nulls.
+ *
+ * @tparam method Whether to use stable sort
+ * @param input The strings column to sort
+ * @param indices The output sorted indices
+ * @param ascending The sort order
+ * @param stream The CUDA stream to use
+ */
+template <sort_method method>
+void sorted_order_strings(column_view const& input,
+                          mutable_column_view& indices,
+                          bool ascending,
+                          cuda::stream_ref stream);
+
+}  // namespace detail
+}  // namespace cudf

cpp/src/sort/sorted_order_strings.cu

@@ -0,0 +1,86 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include "sort_strings.hpp"
+
+#include <cudf/column/column_device_view.cuh>
+#include <cudf/column/column_view.hpp>
+#include <cudf/strings/string_view.cuh>
+#include <cudf/utilities/error.hpp>
+#include <cudf/utilities/memory_resource.hpp>
+
+#include <cub/device/device_merge_sort.cuh>
+#include <cuda/iterator>
+#include <cuda/std/execution>
+#include <cuda/stream>
+
+namespace cudf {
+namespace detail {
+namespace {
+
+/**
+ * @brief Compares two rows of the input strings column by index
+ *
+ * The input column is known to have no nulls so the strings can be
+ * dereferenced and compared without any null checking.
+ */
+struct strings_comparator_fn {
+  column_device_view const d_strings;
+  bool ascending;
+  __device__ bool operator()(size_type lhs, size_type rhs) const
+  {
+    auto const lhs_str = d_strings.element<string_view>(lhs);
+    auto const rhs_str = d_strings.element<string_view>(rhs);
+    return ascending ? (lhs_str < rhs_str) : (rhs_str < lhs_str);
+  }
+};
+
+}  // namespace
+
+bool is_strings_sortable(column_view const& column)
+{
+  return column.type().id() == type_id::STRING && !column.has_nulls() && column.size() > 0;
+}
+
+template <sort_method method>
+void sorted_order_strings(column_view const& input,
+                          mutable_column_view& indices,
+                          bool ascending,
+                          cuda::stream_ref stream)
+{
+  auto const n       = input.size();
+  auto const d_input = column_device_view::create(input, stream);
+
+  // only the indices are sorted; the comparator dereferences the strings on demand so no
+  // string_view array is ever materialized
+  auto const in_keys  = cuda::counting_iterator<size_type>{0};
+  auto const out_keys = indices.begin<size_type>();
+
+  auto const comp = strings_comparator_fn{*d_input, ascending};
+  // the environment provides cub with the stream and the memory resource it uses for
+  // its temporary storage
+  auto const mr_env = cuda::std::execution::prop{cuda::mr::get_memory_resource_t{},
+                                                 cudf::get_current_device_resource_ref()};
+  auto const env    = cuda::std::execution::env{cuda::stream_ref{stream.get()}, mr_env};
+  // Compiling the cub sort APIs is expensive so use a constexpr condition
+  // to only compile the one that is needed.
+  if constexpr (method == sort_method::STABLE) {
+    CUDF_CUDA_TRY(cub::DeviceMergeSort::StableSortKeysCopy(in_keys, out_keys, n, comp, env));
+  } else {
+    CUDF_CUDA_TRY(cub::DeviceMergeSort::SortKeysCopy(in_keys, out_keys, n, comp, env));
+  }
+}
+
+template void sorted_order_strings<sort_method::STABLE>(column_view const&,
+                                                        mutable_column_view&,
+                                                        bool,
+                                                        cuda::stream_ref);
+template void sorted_order_strings<sort_method::UNSTABLE>(column_view const&,
+                                                          mutable_column_view&,
+                                                          bool,
+                                                          cuda::stream_ref);
+
+}  // namespace detail
+}  // namespace cudf

cpp/src/sort/stable_sort_column.cu

@@ -5,6 +5,7 @@
 
 #include "sort_column_impl.cuh"
 #include "sort_radix.hpp"
+#include "sort_strings.hpp"
 
 #include <cudf/column/column_factories.hpp>
 #include <cudf/column/column_view.hpp>
@@ -31,6 +32,9 @@ std::unique_ptr<column> sorted_order<sort_method::STABLE>(column_view const& inp
   mutable_column_view indices_view = sorted_indices->mutable_view();
   if (is_radix_sortable(input)) {
     sorted_order_radix(input, indices_view, column_order == order::ASCENDING, stream);
+  } else if (is_strings_sortable(input)) {
+    sorted_order_strings<sort_method::STABLE>(
+      input, indices_view, column_order == order::ASCENDING, stream);
   } else {
     cudf::type_dispatcher<dispatch_storage_type>(input.type(),
                                                  column_sorted_order_fn<sort_method::STABLE>{},

cpp/tests/sort/stable_sort_tests.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -288,3 +288,24 @@ TEST_F(StableSortDouble, InfinityAndNaN)
   auto results = stable_sorted_order(cudf::table_view({input}));
   CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(results->view(), expected);
 }
+
+struct StableSortStrings : public cudf::test::BaseFixture {};
+
+TEST_F(StableSortStrings, SingleColumnNoNull)
+{
+  // This test exercises the "fast-path" single strings column sort.
+  // Equivalent strings must retain their relative ordering.
+  //                                              0    1    2    3    4    5   6    7
+  cudf::test::strings_column_wrapper col({"cc", "aa", "bb", "aa", "cc", "", "bb", "aa"});
+  auto const input = cudf::table_view({col});
+
+  auto const expected_asc =
+    cudf::test::fixed_width_column_wrapper<cudf::size_type>({5, 1, 3, 7, 2, 6, 0, 4});
+  auto const results_asc = cudf::stable_sorted_order(input, {cudf::order::ASCENDING});
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(results_asc->view(), expected_asc);
+
+  auto const expected_desc =
+    cudf::test::fixed_width_column_wrapper<cudf::size_type>({0, 4, 2, 6, 1, 3, 7, 5});
+  auto const results_desc = cudf::stable_sorted_order(input, {cudf::order::DESCENDING});
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(results_desc->view(), expected_desc);
+}