From 57bb25fa30b6d1af17c9d26eafa61850dcd62c2a Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Sun, 19 Jul 2026 06:57:00 +0000
Subject: [PATCH 1/4] Reduce compute_global_memory_aggs_null build time

---
 cpp/CMakeLists.txt                            |   6 +
 .../detail/aggregation/device_aggregators.cuh |  20 +-
 .../hash/compute_global_memory_aggs.hpp       |   4 +-
 .../hash/compute_global_memory_aggs_null.cu   |  31 ++-
 .../hash/compute_global_memory_aggs_null.hpp  |  40 ++++
 .../compute_global_memory_aggs_null_dense.cu  |  74 ++++++
 ...lobal_memory_aggs_null_dense_dictionary.cu |  21 ++
 ...l_memory_aggs_null_dense_non_dictionary.cu |  21 ++
 ...ompute_global_memory_aggs_null_kernels.cuh |  71 ++++++
 ...ompute_global_memory_aggs_null_kernels.hpp |  52 ++++
 .../compute_global_memory_aggs_null_sparse.cu |  66 ++++++
 ...obal_memory_aggs_null_sparse_dictionary.cu |  22 ++
 ..._memory_aggs_null_sparse_non_dictionary.cu |  22 ++
 cpp/src/groupby/hash/single_pass_functors.cuh | 223 ++++++++++++++++--
 14 files changed, 644 insertions(+), 29 deletions(-)
 create mode 100644 cpp/src/groupby/hash/compute_global_memory_aggs_null.hpp
 create mode 100644 cpp/src/groupby/hash/compute_global_memory_aggs_null_dense.cu
 create mode 100644 cpp/src/groupby/hash/compute_global_memory_aggs_null_dense_dictionary.cu
 create mode 100644 cpp/src/groupby/hash/compute_global_memory_aggs_null_dense_non_dictionary.cu
 create mode 100644 cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.cuh
 create mode 100644 cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.hpp
 create mode 100644 cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse.cu
 create mode 100644 cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse_dictionary.cu
 create mode 100644 cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse_non_dictionary.cu

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 0899c3cec002..9dc520061d5c 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -662,6 +662,12 @@ add_library(
   src/groupby/groupby.cu
   src/groupby/hash/compute_global_memory_aggs.cu
   src/groupby/hash/compute_global_memory_aggs_null.cu
+  src/groupby/hash/compute_global_memory_aggs_null_dense.cu
+  src/groupby/hash/compute_global_memory_aggs_null_dense_dictionary.cu
+  src/groupby/hash/compute_global_memory_aggs_null_dense_non_dictionary.cu
+  src/groupby/hash/compute_global_memory_aggs_null_sparse.cu
+  src/groupby/hash/compute_global_memory_aggs_null_sparse_dictionary.cu
+  src/groupby/hash/compute_global_memory_aggs_null_sparse_non_dictionary.cu
   src/groupby/hash/compute_groupby.cu
   src/groupby/hash/compute_mapping_indices.cu
   src/groupby/hash/compute_mapping_indices_null.cu
diff --git a/cpp/include/cudf/detail/aggregation/device_aggregators.cuh b/cpp/include/cudf/detail/aggregation/device_aggregators.cuh
index 3e919f76ccaf..a12ecdf50224 100644
--- a/cpp/include/cudf/detail/aggregation/device_aggregators.cuh
+++ b/cpp/include/cudf/detail/aggregation/device_aggregators.cuh
@@ -166,9 +166,11 @@ struct update_target_element<Source, aggregation::SUM_OVERFLOW> {
  * SFINAE is used to prevent recursion for dictionary type. Dictionary keys cannot be a
  * dictionary.
  *
+ * @tparam k Aggregation to perform
  */
+template <aggregation::Kind k>
 struct update_target_from_dictionary {
-  template <typename Source, aggregation::Kind k>
+  template <typename Source>
   __device__ void operator()(mutable_column_device_view target,
                              size_type target_index,
                              column_device_view source,
@@ -177,7 +179,7 @@ struct update_target_from_dictionary {
   {
     update_target_element<Source, k>{}(target, target_index, source, source_index);
   }
-  template <typename Source, aggregation::Kind k>
+  template <typename Source>
   __device__ void operator()(mutable_column_device_view,
                              size_type,
                              column_device_view,
@@ -207,14 +209,12 @@ struct update_target_element<dictionary32, k> {
                              column_device_view source,
                              size_type source_index) const noexcept
   {
-    dispatch_type_and_aggregation(
-      source.child(cudf::dictionary_column_view::keys_column_index).type(),
-      k,
-      update_target_from_dictionary{},
-      target,
-      target_index,
-      source.child(cudf::dictionary_column_view::keys_column_index),
-      static_cast<cudf::size_type>(source.element<dictionary32>(source_index)));
+    type_dispatcher(source.child(cudf::dictionary_column_view::keys_column_index).type(),
+                    update_target_from_dictionary<k>{},
+                    target,
+                    target_index,
+                    source.child(cudf::dictionary_column_view::keys_column_index),
+                    static_cast<cudf::size_type>(source.element<dictionary32>(source_index)));
   }
 };
 
diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs.hpp b/cpp/src/groupby/hash/compute_global_memory_aggs.hpp
index ef9d730d7118..1990d0cc2d99 100644
--- a/cpp/src/groupby/hash/compute_global_memory_aggs.hpp
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs.hpp
@@ -5,11 +5,13 @@
 #pragma once
 
 #include <cudf/aggregation.hpp>
-#include <cudf/table/table_device_view.cuh>
+#include <cudf/table/table.hpp>
+#include <cudf/table/table_view.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/device_uvector.hpp>
 #include <rmm/resource_ref.hpp>
 
 #include <cstdint>
diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null.cu b/cpp/src/groupby/hash/compute_global_memory_aggs_null.cu
index d50c7d302710..209c0a0ce4f9 100644
--- a/cpp/src/groupby/hash/compute_global_memory_aggs_null.cu
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null.cu
@@ -3,13 +3,19 @@
  * SPDX-License-Identifier: Apache-2.0
  */
 
-#include "compute_global_memory_aggs.cuh"
+#include "compute_global_memory_aggs.hpp"
+#include "compute_global_memory_aggs_null.hpp"
+#include "helpers.cuh"
 
+#include <cstdint>
+#include <memory>
 #include <span>
+#include <utility>
 
 namespace cudf::groupby::detail::hash {
 
-template std::pair<std::unique_ptr<table>, rmm::device_uvector<size_type>>
+template <>
+std::pair<std::unique_ptr<table>, rmm::device_uvector<size_type>>
 compute_global_memory_aggs<nullable_global_set_t>(bitmask_type const* row_bitmask,
                                                   table_view const& values,
                                                   nullable_global_set_t const& key_set,
@@ -17,6 +23,25 @@ compute_global_memory_aggs<nullable_global_set_t>(bitmask_type const* row_bitmas
                                                   device_span<aggregation::Kind const> d_agg_kinds,
                                                   std::span<int8_t const> is_agg_intermediate,
                                                   rmm::cuda_stream_view stream,
-                                                  rmm::device_async_resource_ref mr);
+                                                  rmm::device_async_resource_ref mr)
+{
+  return h_agg_kinds.size() > GROUPBY_DENSE_OUTPUT_THRESHOLD
+           ? compute_global_memory_aggs_null_dense(row_bitmask,
+                                                   values,
+                                                   key_set,
+                                                   h_agg_kinds,
+                                                   d_agg_kinds,
+                                                   is_agg_intermediate,
+                                                   stream,
+                                                   mr)
+           : compute_global_memory_aggs_null_sparse(row_bitmask,
+                                                    values,
+                                                    key_set,
+                                                    h_agg_kinds,
+                                                    d_agg_kinds,
+                                                    is_agg_intermediate,
+                                                    stream,
+                                                    mr);
+}
 
 }  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null.hpp b/cpp/src/groupby/hash/compute_global_memory_aggs_null.hpp
new file mode 100644
index 000000000000..11d93da5a034
--- /dev/null
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null.hpp
@@ -0,0 +1,40 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#pragma once
+
+#include "compute_global_memory_aggs.hpp"
+#include "helpers.cuh"
+
+#include <cstdint>
+#include <memory>
+#include <span>
+#include <utility>
+
+namespace cudf::groupby::detail::hash {
+
+std::pair<std::unique_ptr<table>, rmm::device_uvector<size_type>>
+compute_global_memory_aggs_null_dense(
+  bitmask_type const* row_bitmask,
+  table_view const& values,
+  nullable_global_set_t const& key_set,
+  host_span<aggregation::Kind const> h_agg_kinds,
+  device_span<aggregation::Kind const> d_agg_kinds,
+  std::span<int8_t const> is_agg_intermediate,
+  rmm::cuda_stream_view stream,
+  rmm::device_async_resource_ref mr);
+
+std::pair<std::unique_ptr<table>, rmm::device_uvector<size_type>>
+compute_global_memory_aggs_null_sparse(
+  bitmask_type const* row_bitmask,
+  table_view const& values,
+  nullable_global_set_t const& key_set,
+  host_span<aggregation::Kind const> h_agg_kinds,
+  device_span<aggregation::Kind const> d_agg_kinds,
+  std::span<int8_t const> is_agg_intermediate,
+  rmm::cuda_stream_view stream,
+  rmm::device_async_resource_ref mr);
+
+}  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null_dense.cu b/cpp/src/groupby/hash/compute_global_memory_aggs_null_dense.cu
new file mode 100644
index 000000000000..cf2f9a37b165
--- /dev/null
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null_dense.cu
@@ -0,0 +1,74 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include "compute_global_memory_aggs.cuh"
+#include "compute_global_memory_aggs_null.hpp"
+#include "compute_global_memory_aggs_null_kernels.hpp"
+
+#include <cudf/column/column_view.hpp>
+#include <cudf/table/table_device_view.cuh>
+#include <cudf/utilities/memory_resource.hpp>
+#include <cudf/utilities/traits.hpp>
+
+#include <algorithm>
+#include <cstdint>
+#include <memory>
+#include <span>
+#include <utility>
+
+namespace cudf::groupby::detail::hash {
+
+std::pair<std::unique_ptr<table>, rmm::device_uvector<size_type>>
+compute_global_memory_aggs_null_dense(bitmask_type const* row_bitmask,
+                                      table_view const& values,
+                                      nullable_global_set_t const& key_set,
+                                      host_span<aggregation::Kind const> h_agg_kinds,
+                                      device_span<aggregation::Kind const> d_agg_kinds,
+                                      std::span<int8_t const> is_agg_intermediate,
+                                      rmm::cuda_stream_view stream,
+                                      rmm::device_async_resource_ref mr)
+{
+  auto const num_rows                = values.num_rows();
+  auto [unique_keys, target_indices] = [&] {
+    auto matching_keys =
+      compute_matching_keys(row_bitmask, key_set.ref(cuco::op::insert_and_find), num_rows, stream);
+    auto unique_keys       = extract_populated_keys(key_set, num_rows, stream, mr);
+    auto key_transform_map = compute_key_transform_map(
+      num_rows, unique_keys, stream, cudf::get_current_device_resource_ref());
+    auto target_indices = compute_target_indices(
+      matching_keys, key_transform_map, stream, cudf::get_current_device_resource_ref());
+    return std::pair{std::move(unique_keys), std::move(target_indices)};
+  }();
+
+  auto const d_values  = table_device_view::create(values, stream);
+  auto agg_results     = create_results_table(static_cast<size_type>(unique_keys.size()),
+                                          values,
+                                          h_agg_kinds,
+                                          is_agg_intermediate,
+                                          stream,
+                                          mr);
+  auto d_results       = mutable_table_device_view::create(*agg_results, stream);
+  auto const num_items = num_rows * static_cast<int64_t>(h_agg_kinds.size());
+
+  auto const has_dictionary = std::any_of(
+    values.begin(), values.end(), [](column_view const& col) { return is_dictionary(col.type()); });
+  auto const has_non_dictionary =
+    std::any_of(values.begin(), values.end(), [](column_view const& col) {
+      return not is_dictionary(col.type());
+    });
+
+  if (has_non_dictionary) {
+    launch_null_dense_non_dictionary(
+      target_indices.data(), d_agg_kinds.data(), *d_values, *d_results, num_items, stream);
+  }
+  if (has_dictionary) {
+    launch_null_dense_dictionary(
+      target_indices.data(), d_agg_kinds.data(), *d_values, *d_results, num_items, stream);
+  }
+
+  return {std::move(agg_results), std::move(unique_keys)};
+}
+
+}  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null_dense_dictionary.cu b/cpp/src/groupby/hash/compute_global_memory_aggs_null_dense_dictionary.cu
new file mode 100644
index 000000000000..bf151c48d8de
--- /dev/null
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null_dense_dictionary.cu
@@ -0,0 +1,21 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include "compute_global_memory_aggs_null_kernels.cuh"
+
+namespace cudf::groupby::detail::hash {
+
+void launch_null_dense_dictionary(size_type const* target_indices,
+                                  aggregation::Kind const* aggs,
+                                  table_device_view const& input_values,
+                                  mutable_table_device_view const& output_values,
+                                  int64_t num_items,
+                                  rmm::cuda_stream_view stream)
+{
+  launch_null_dense_filtered<true>(
+    target_indices, aggs, input_values, output_values, num_items, stream);
+}
+
+}  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null_dense_non_dictionary.cu b/cpp/src/groupby/hash/compute_global_memory_aggs_null_dense_non_dictionary.cu
new file mode 100644
index 000000000000..7649cd2ef839
--- /dev/null
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null_dense_non_dictionary.cu
@@ -0,0 +1,21 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include "compute_global_memory_aggs_null_kernels.cuh"
+
+namespace cudf::groupby::detail::hash {
+
+void launch_null_dense_non_dictionary(size_type const* target_indices,
+                                      aggregation::Kind const* aggs,
+                                      table_device_view const& input_values,
+                                      mutable_table_device_view const& output_values,
+                                      int64_t num_items,
+                                      rmm::cuda_stream_view stream)
+{
+  launch_null_dense_filtered<false>(
+    target_indices, aggs, input_values, output_values, num_items, stream);
+}
+
+}  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.cuh b/cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.cuh
new file mode 100644
index 000000000000..b036d46dc407
--- /dev/null
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.cuh
@@ -0,0 +1,71 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#pragma once
+
+#include "compute_global_memory_aggs_null_kernels.hpp"
+#include "single_pass_functors.cuh"
+
+#include <cudf/detail/utilities/grid_1d.cuh>
+#include <cudf/utilities/error.hpp>
+
+#include <cuda_runtime_api.h>
+
+namespace cudf::groupby::detail::hash {
+
+template <typename Index, typename Function>
+CUDF_KERNEL void filtered_single_pass_aggs_kernel(Index num_items, Function fn)
+{
+  auto const idx = static_cast<Index>(cudf::detail::grid_1d::global_thread_id());
+  if (idx < num_items) { fn(idx); }
+}
+
+template <typename Index, typename Function>
+void launch_filtered_single_pass_aggs(Index num_items, Function fn, rmm::cuda_stream_view stream)
+{
+  if (num_items == 0) { return; }
+
+  // Match the launch geometry used by the original Thrust kernel. A smaller block size causes a
+  // significant runtime regression in the hash table operations.
+  constexpr auto block_size = 256;
+  cudf::detail::grid_1d config{num_items, block_size};
+  filtered_single_pass_aggs_kernel<<<config.num_blocks,
+                                     config.num_threads_per_block,
+                                     0,
+                                     stream.value()>>>(num_items, fn);
+  CUDF_CUDA_TRY(cudaGetLastError());
+}
+
+template <bool IsDictionary>
+void launch_null_sparse_filtered(nullable_insert_and_find_ref set_ref,
+                                 bitmask_type const* row_bitmask,
+                                 aggregation::Kind const* aggs,
+                                 table_device_view const& input_values,
+                                 mutable_table_device_view const& output_values,
+                                 size_type num_rows,
+                                 rmm::cuda_stream_view stream)
+{
+  launch_filtered_single_pass_aggs(
+    num_rows,
+    compute_filtered_single_pass_aggs_sparse_output_fn<IsDictionary, nullable_insert_and_find_ref>{
+      set_ref, row_bitmask, aggs, input_values, output_values},
+    stream);
+}
+
+template <bool IsDictionary>
+void launch_null_dense_filtered(size_type const* target_indices,
+                                aggregation::Kind const* aggs,
+                                table_device_view const& input_values,
+                                mutable_table_device_view const& output_values,
+                                int64_t num_items,
+                                rmm::cuda_stream_view stream)
+{
+  launch_filtered_single_pass_aggs(num_items,
+                                   compute_filtered_single_pass_aggs_dense_output_fn<IsDictionary>{
+                                     target_indices, aggs, input_values, output_values},
+                                   stream);
+}
+
+}  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.hpp b/cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.hpp
new file mode 100644
index 000000000000..8b4eb0dca9d2
--- /dev/null
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.hpp
@@ -0,0 +1,52 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#pragma once
+
+#include "helpers.cuh"
+
+#include <cudf/aggregation.hpp>
+#include <cudf/table/table_device_view.cuh>
+#include <cudf/types.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+
+#include <cstdint>
+
+namespace cudf::groupby::detail::hash {
+
+using nullable_insert_and_find_ref = nullable_hash_set_ref_t<cuco::insert_and_find_tag>;
+
+void launch_null_sparse_non_dictionary(nullable_insert_and_find_ref set_ref,
+                                       bitmask_type const* row_bitmask,
+                                       aggregation::Kind const* aggs,
+                                       table_device_view const& input_values,
+                                       mutable_table_device_view const& output_values,
+                                       size_type num_rows,
+                                       rmm::cuda_stream_view stream);
+
+void launch_null_sparse_dictionary(nullable_insert_and_find_ref set_ref,
+                                   bitmask_type const* row_bitmask,
+                                   aggregation::Kind const* aggs,
+                                   table_device_view const& input_values,
+                                   mutable_table_device_view const& output_values,
+                                   size_type num_rows,
+                                   rmm::cuda_stream_view stream);
+
+void launch_null_dense_non_dictionary(size_type const* target_indices,
+                                      aggregation::Kind const* aggs,
+                                      table_device_view const& input_values,
+                                      mutable_table_device_view const& output_values,
+                                      int64_t num_items,
+                                      rmm::cuda_stream_view stream);
+
+void launch_null_dense_dictionary(size_type const* target_indices,
+                                  aggregation::Kind const* aggs,
+                                  table_device_view const& input_values,
+                                  mutable_table_device_view const& output_values,
+                                  int64_t num_items,
+                                  rmm::cuda_stream_view stream);
+
+}  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse.cu b/cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse.cu
new file mode 100644
index 000000000000..e80579a27b1e
--- /dev/null
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse.cu
@@ -0,0 +1,66 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include "compute_global_memory_aggs_null.hpp"
+#include "compute_global_memory_aggs_null_kernels.hpp"
+#include "output_utils.hpp"
+
+#include <cudf/column/column_view.hpp>
+#include <cudf/detail/gather.hpp>
+#include <cudf/table/table_device_view.cuh>
+#include <cudf/utilities/traits.hpp>
+
+#include <algorithm>
+#include <cstdint>
+#include <memory>
+#include <span>
+#include <utility>
+
+namespace cudf::groupby::detail::hash {
+
+std::pair<std::unique_ptr<table>, rmm::device_uvector<size_type>>
+compute_global_memory_aggs_null_sparse(bitmask_type const* row_bitmask,
+                                       table_view const& values,
+                                       nullable_global_set_t const& key_set,
+                                       host_span<aggregation::Kind const> h_agg_kinds,
+                                       device_span<aggregation::Kind const> d_agg_kinds,
+                                       std::span<int8_t const> is_agg_intermediate,
+                                       rmm::cuda_stream_view stream,
+                                       rmm::device_async_resource_ref mr)
+{
+  auto const num_rows = values.num_rows();
+  auto const d_values = table_device_view::create(values, stream);
+  auto agg_results =
+    create_results_table(num_rows, values, h_agg_kinds, is_agg_intermediate, stream, mr);
+  auto d_results = mutable_table_device_view::create(*agg_results, stream);
+
+  auto const has_dictionary = std::any_of(
+    values.begin(), values.end(), [](column_view const& col) { return is_dictionary(col.type()); });
+  auto const has_non_dictionary =
+    std::any_of(values.begin(), values.end(), [](column_view const& col) {
+      return not is_dictionary(col.type());
+    });
+  auto const set_ref = key_set.ref(cuco::op::insert_and_find);
+
+  if (has_non_dictionary) {
+    launch_null_sparse_non_dictionary(
+      set_ref, row_bitmask, d_agg_kinds.data(), *d_values, *d_results, num_rows, stream);
+  }
+  if (has_dictionary) {
+    launch_null_sparse_dictionary(
+      set_ref, row_bitmask, d_agg_kinds.data(), *d_values, *d_results, num_rows, stream);
+  }
+
+  auto unique_keys   = extract_populated_keys(key_set, num_rows, stream, mr);
+  auto dense_results = cudf::detail::gather(agg_results->view(),
+                                            unique_keys,
+                                            out_of_bounds_policy::DONT_CHECK,
+                                            cudf::negative_index_policy::NOT_ALLOWED,
+                                            stream,
+                                            mr);
+  return {std::move(dense_results), std::move(unique_keys)};
+}
+
+}  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse_dictionary.cu b/cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse_dictionary.cu
new file mode 100644
index 000000000000..d9289a099a1c
--- /dev/null
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse_dictionary.cu
@@ -0,0 +1,22 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include "compute_global_memory_aggs_null_kernels.cuh"
+
+namespace cudf::groupby::detail::hash {
+
+void launch_null_sparse_dictionary(nullable_insert_and_find_ref set_ref,
+                                   bitmask_type const* row_bitmask,
+                                   aggregation::Kind const* aggs,
+                                   table_device_view const& input_values,
+                                   mutable_table_device_view const& output_values,
+                                   size_type num_rows,
+                                   rmm::cuda_stream_view stream)
+{
+  launch_null_sparse_filtered<true>(
+    set_ref, row_bitmask, aggs, input_values, output_values, num_rows, stream);
+}
+
+}  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse_non_dictionary.cu b/cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse_non_dictionary.cu
new file mode 100644
index 000000000000..a06f056075d2
--- /dev/null
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null_sparse_non_dictionary.cu
@@ -0,0 +1,22 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include "compute_global_memory_aggs_null_kernels.cuh"
+
+namespace cudf::groupby::detail::hash {
+
+void launch_null_sparse_non_dictionary(nullable_insert_and_find_ref set_ref,
+                                       bitmask_type const* row_bitmask,
+                                       aggregation::Kind const* aggs,
+                                       table_device_view const& input_values,
+                                       mutable_table_device_view const& output_values,
+                                       size_type num_rows,
+                                       rmm::cuda_stream_view stream)
+{
+  launch_null_sparse_filtered<false>(
+    set_ref, row_bitmask, aggs, input_values, output_values, num_rows, stream);
+}
+
+}  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/groupby/hash/single_pass_functors.cuh b/cpp/src/groupby/hash/single_pass_functors.cuh
index 165b7fa2bc41..c22cef57a4cc 100644
--- a/cpp/src/groupby/hash/single_pass_functors.cuh
+++ b/cpp/src/groupby/hash/single_pass_functors.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -9,10 +9,115 @@
 #include <cudf/detail/aggregation/aggregation.cuh>
 #include <cudf/detail/aggregation/device_aggregators.cuh>
 #include <cudf/detail/utilities/assert.cuh>
+#include <cudf/table/table_device_view.cuh>
+#include <cudf/types.hpp>
+#include <cudf/utilities/bit.hpp>
+#include <cudf/utilities/type_dispatcher.hpp>
 
 #include <cuda/std/cstddef>
+#include <cuda/std/type_traits>
+#include <cuda/std/utility>
+
+#include <cstdint>
 
 namespace cudf::groupby::detail::hash {
+
+// Compound hash aggregations are decomposed into these simple aggregations before these kernels
+// are launched.
+template <typename F, typename... Ts>
+__device__ auto dispatch_single_pass_aggregation(cudf::aggregation::Kind kind, F&& f, Ts&&... args)
+{
+  switch (kind) {
+    case cudf::aggregation::SUM:
+      return f.template operator()<cudf::aggregation::SUM>(cuda::std::forward<Ts>(args)...);
+    case cudf::aggregation::SUM_OVERFLOW:
+      return f.template operator()<cudf::aggregation::SUM_OVERFLOW>(
+        cuda::std::forward<Ts>(args)...);
+    case cudf::aggregation::PRODUCT:
+      return f.template operator()<cudf::aggregation::PRODUCT>(cuda::std::forward<Ts>(args)...);
+    case cudf::aggregation::MIN:
+      return f.template operator()<cudf::aggregation::MIN>(cuda::std::forward<Ts>(args)...);
+    case cudf::aggregation::MAX:
+      return f.template operator()<cudf::aggregation::MAX>(cuda::std::forward<Ts>(args)...);
+    case cudf::aggregation::COUNT_VALID:
+      return f.template operator()<cudf::aggregation::COUNT_VALID>(cuda::std::forward<Ts>(args)...);
+    case cudf::aggregation::COUNT_ALL:
+      return f.template operator()<cudf::aggregation::COUNT_ALL>(cuda::std::forward<Ts>(args)...);
+    case cudf::aggregation::SUM_OF_SQUARES:
+      return f.template operator()<cudf::aggregation::SUM_OF_SQUARES>(
+        cuda::std::forward<Ts>(args)...);
+    case cudf::aggregation::ARGMAX:
+      return f.template operator()<cudf::aggregation::ARGMAX>(cuda::std::forward<Ts>(args)...);
+    case cudf::aggregation::ARGMIN:
+      return f.template operator()<cudf::aggregation::ARGMIN>(cuda::std::forward<Ts>(args)...);
+    default: CUDF_UNREACHABLE("Unsupported single-pass aggregation.");
+  }
+}
+
+struct unsupported_hash_value_type {};
+
+template <cudf::type_id Id>
+struct dispatch_hash_value_type {
+  // Nested value columns are rejected by can_use_hash_groupby before these kernels are launched.
+  using type = cuda::std::conditional_t<Id == cudf::type_id::LIST or Id == cudf::type_id::STRUCT,
+                                        unsupported_hash_value_type,
+                                        cudf::id_to_type<Id>>;
+};
+
+template <cudf::type_id Id>
+struct dispatch_non_dictionary_hash_value_type {
+  using type = cuda::std::conditional_t<Id == cudf::type_id::DICTIONARY32 or
+                                          Id == cudf::type_id::LIST or Id == cudf::type_id::STRUCT,
+                                        unsupported_hash_value_type,
+                                        cudf::id_to_type<Id>>;
+};
+
+template <cudf::type_id Id>
+struct dispatch_dictionary_hash_value_type {
+  using type = cuda::std::conditional_t<Id == cudf::type_id::DICTIONARY32,
+                                        cudf::id_to_type<Id>,
+                                        unsupported_hash_value_type>;
+};
+
+template <typename Element>
+struct dispatch_single_pass_aggregation_fn {
+  template <cudf::aggregation::Kind kind, typename F, typename... Ts>
+  __device__ auto operator()(F&& f, Ts&&... args) const
+  {
+    return f.template operator()<Element, kind>(cuda::std::forward<Ts>(args)...);
+  }
+};
+
+struct dispatch_single_pass_source_fn {
+  template <typename Element, typename F, typename... Ts>
+  __device__ auto operator()(cudf::aggregation::Kind kind, F&& f, Ts&&... args) const
+  {
+    if constexpr (cuda::std::is_same_v<Element, unsupported_hash_value_type>) {
+      CUDF_UNREACHABLE("Unsupported hash groupby value type.");
+    } else {
+      return dispatch_single_pass_aggregation(kind,
+                                              dispatch_single_pass_aggregation_fn<Element>{},
+                                              cuda::std::forward<F>(f),
+                                              cuda::std::forward<Ts>(args)...);
+    }
+  }
+};
+
+template <template <cudf::type_id> class DispatchMap = dispatch_hash_value_type,
+          typename F,
+          typename... Ts>
+__device__ auto dispatch_single_pass_type_and_aggregation(cudf::data_type type,
+                                                          cudf::aggregation::Kind kind,
+                                                          F&& f,
+                                                          Ts&&... args)
+{
+  return cudf::type_dispatcher<DispatchMap>(type,
+                                            dispatch_single_pass_source_fn{},
+                                            kind,
+                                            cuda::std::forward<F>(f),
+                                            cuda::std::forward<Ts>(args)...);
+}
+
 /// Functor used by type dispatcher returning the size of the underlying C++ type
 struct size_of_functor {
   template <typename T>
@@ -106,13 +211,13 @@ struct compute_single_pass_aggs_sparse_output_fn : compute_single_pass_aggs_base
     for (size_type col_idx = 0; col_idx < input_values.num_columns(); ++col_idx) {
       auto const& source_col = input_values.column(col_idx);
       auto const& target_col = output_values.column(col_idx);
-      dispatch_type_and_aggregation(source_col.type(),
-                                    aggs[col_idx],
-                                    cudf::detail::element_aggregator{},
-                                    target_col,
-                                    target_row_idx,
-                                    source_col,
-                                    idx);
+      dispatch_single_pass_type_and_aggregation(source_col.type(),
+                                                aggs[col_idx],
+                                                cudf::detail::element_aggregator{},
+                                                target_col,
+                                                target_row_idx,
+                                                source_col,
+                                                idx);
     }
   }
 };
@@ -145,15 +250,103 @@ struct compute_single_pass_aggs_dense_output_fn : compute_single_pass_aggs_base_
       auto const col_idx     = static_cast<size_type>(idx / num_rows);
       auto const& source_col = input_values.column(col_idx);
       auto const& target_col = output_values.column(col_idx);
-      dispatch_type_and_aggregation(source_col.type(),
-                                    aggs[col_idx],
-                                    cudf::detail::element_aggregator{},
-                                    target_col,
-                                    target_row_idx,
-                                    source_col,
-                                    source_row_idx);
+      dispatch_single_pass_type_and_aggregation(source_col.type(),
+                                                aggs[col_idx],
+                                                cudf::detail::element_aggregator{},
+                                                target_col,
+                                                target_row_idx,
+                                                source_col,
+                                                source_row_idx);
+    }
+  }
+};
+
+template <bool IsDictionary>
+__device__ void aggregate_filtered_single_pass_value(cudf::aggregation::Kind kind,
+                                                     mutable_column_device_view target,
+                                                     size_type target_index,
+                                                     column_device_view source,
+                                                     size_type source_index)
+{
+  if constexpr (IsDictionary) {
+    dispatch_single_pass_type_and_aggregation<dispatch_dictionary_hash_value_type>(
+      source.type(),
+      kind,
+      cudf::detail::element_aggregator{},
+      target,
+      target_index,
+      source,
+      source_index);
+  } else {
+    dispatch_single_pass_type_and_aggregation<dispatch_non_dictionary_hash_value_type>(
+      source.type(),
+      kind,
+      cudf::detail::element_aggregator{},
+      target,
+      target_index,
+      source,
+      source_index);
+  }
+}
+
+template <bool IsDictionary, typename SetRef>
+struct compute_filtered_single_pass_aggs_sparse_output_fn : compute_single_pass_aggs_base_fn {
+  SetRef set_ref;
+  bitmask_type const* row_bitmask;
+
+  compute_filtered_single_pass_aggs_sparse_output_fn(SetRef set_ref,
+                                                     bitmask_type const* row_bitmask,
+                                                     aggregation::Kind const* aggs,
+                                                     table_device_view const& input_values,
+                                                     mutable_table_device_view const& output_values)
+    : compute_single_pass_aggs_base_fn(aggs, input_values, output_values),
+      set_ref{set_ref},
+      row_bitmask{row_bitmask}
+  {
+  }
+
+  __device__ void operator()(size_type idx)
+  {
+    if (row_bitmask && !cudf::bit_is_set(row_bitmask, idx)) { return; }
+    auto const target_row_idx = *set_ref.insert_and_find(idx).first;
+
+    for (size_type col_idx = 0; col_idx < input_values.num_columns(); ++col_idx) {
+      auto const& source_col = input_values.column(col_idx);
+      if ((source_col.type().id() == type_id::DICTIONARY32) != IsDictionary) { continue; }
+
+      aggregate_filtered_single_pass_value<IsDictionary>(
+        aggs[col_idx], output_values.column(col_idx), target_row_idx, source_col, idx);
     }
   }
 };
 
+template <bool IsDictionary>
+struct compute_filtered_single_pass_aggs_dense_output_fn : compute_single_pass_aggs_base_fn {
+  size_type const* target_indices;
+
+  compute_filtered_single_pass_aggs_dense_output_fn(size_type const* target_indices,
+                                                    aggregation::Kind const* aggs,
+                                                    table_device_view const& input_values,
+                                                    mutable_table_device_view const& output_values)
+    : compute_single_pass_aggs_base_fn(aggs, input_values, output_values),
+      target_indices(target_indices)
+  {
+  }
+
+  __device__ void operator()(int64_t idx) const
+  {
+    auto const num_rows       = input_values.num_rows();
+    auto const source_row_idx = static_cast<size_type>(idx % num_rows);
+    auto const target_row_idx = target_indices[source_row_idx];
+    if (target_row_idx == cudf::detail::CUDF_SIZE_TYPE_SENTINEL) { return; }
+
+    auto const col_idx     = static_cast<size_type>(idx / num_rows);
+    auto const& source_col = input_values.column(col_idx);
+    if ((source_col.type().id() == type_id::DICTIONARY32) != IsDictionary) { return; }
+
+    aggregate_filtered_single_pass_value<IsDictionary>(
+      aggs[col_idx], output_values.column(col_idx), target_row_idx, source_col, source_row_idx);
+  }
+};
+
 }  // namespace cudf::groupby::detail::hash

From 5f597b469b1b9d3af1fd2012a689575badb9e610 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Tue, 8 Sep 2026 13:55:09 -0700
Subject: [PATCH 2/4] Declare the nullable compute_global_memory_aggs
 specialization before use

The nullable key set is now an explicit specialization defined in
compute_global_memory_aggs_null.cu, but callers reached it through
compute_single_pass_aggs.cuh, which only saw the primary template
declaration. An explicit specialization must be declared before any use
that would otherwise instantiate the primary template, in every
translation unit; violating that is ill-formed with no diagnostic
required. Declare the specialization in compute_global_memory_aggs_null.hpp
and include that header from the single-pass aggregation template.
---
 .../hash/compute_global_memory_aggs_null.hpp   | 18 ++++++++++++++++++
 .../groupby/hash/compute_single_pass_aggs.cuh  |  1 +
 2 files changed, 19 insertions(+)

diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null.hpp b/cpp/src/groupby/hash/compute_global_memory_aggs_null.hpp
index f528ee091e51..16c3c35264ea 100644
--- a/cpp/src/groupby/hash/compute_global_memory_aggs_null.hpp
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null.hpp
@@ -15,6 +15,24 @@
 
 namespace cudf::groupby::detail::hash {
 
+/*
+ * The nullable key set does not instantiate the primary template in
+ * `compute_global_memory_aggs.cuh`; it is explicitly specialized in
+ * `compute_global_memory_aggs_null.cu` to dispatch to the split implementations below.  An explicit
+ * specialization must be declared before any use that would otherwise instantiate the primary
+ * template, so every caller that may aggregate with `nullable_global_set_t` includes this header.
+ */
+template <>
+std::pair<std::unique_ptr<table>, rmm::device_uvector<size_type>>
+compute_global_memory_aggs<nullable_global_set_t>(bitmask_type const* row_bitmask,
+                                                  table_view const& values,
+                                                  nullable_global_set_t const& key_set,
+                                                  host_span<aggregation::Kind const> h_agg_kinds,
+                                                  device_span<aggregation::Kind const> d_agg_kinds,
+                                                  std::span<int8_t const> is_agg_intermediate,
+                                                  cuda::stream_ref stream,
+                                                  rmm::device_async_resource_ref mr);
+
 std::pair<std::unique_ptr<table>, rmm::device_uvector<size_type>>
 compute_global_memory_aggs_null_dense(bitmask_type const* row_bitmask,
                                       table_view const& values,
diff --git a/cpp/src/groupby/hash/compute_single_pass_aggs.cuh b/cpp/src/groupby/hash/compute_single_pass_aggs.cuh
index d38acb4b6522..e360fe2ea674 100644
--- a/cpp/src/groupby/hash/compute_single_pass_aggs.cuh
+++ b/cpp/src/groupby/hash/compute_single_pass_aggs.cuh
@@ -6,6 +6,7 @@
 #pragma once
 
 #include "compute_global_memory_aggs.hpp"
+#include "compute_global_memory_aggs_null.hpp"
 #include "compute_mapping_indices.hpp"
 #include "compute_shared_memory_aggs.hpp"
 #include "compute_single_pass_aggs.hpp"

From 7e133c3c663679a8be61886bd39038de95e6247b Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Tue, 8 Sep 2026 13:55:14 -0700
Subject: [PATCH 3/4] Bounds-check the thread index before narrowing in the
 filtered aggs kernel

The kernel cast the 64-bit global thread id to the item index type
before comparing it against num_items. For the sparse path the index
type is 32-bit, so thread ids in the final partial block could wrap when
the row count is near the size_type limit. Compare in the 64-bit thread
index space first, then narrow, matching the rest of libcudf.
---
 .../hash/compute_global_memory_aggs_null_kernels.cuh       | 7 +++++--
 1 file changed, 5 insertions(+), 2 deletions(-)

diff --git a/cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.cuh b/cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.cuh
index ec5fd4fa2b5a..20b1af385b44 100644
--- a/cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.cuh
+++ b/cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.cuh
@@ -18,8 +18,11 @@ namespace cudf::groupby::detail::hash {
 template <typename Index, typename Function>
 CUDF_KERNEL void filtered_single_pass_aggs_kernel(Index num_items, Function fn)
 {
-  auto const idx = static_cast<Index>(cudf::detail::grid_1d::global_thread_id());
-  if (idx < num_items) { fn(idx); }
+  // The last block can extend past `num_items`, and for a 32-bit `Index` those thread ids may not
+  // be representable, so compare in the 64-bit thread index space before narrowing.
+  auto const tid = cudf::detail::grid_1d::global_thread_id();
+  if (tid >= static_cast<cudf::thread_index_type>(num_items)) { return; }
+  fn(static_cast<Index>(tid));
 }
 
 template <typename Index, typename Function>

From 3d5da50c872dbfa42112aa020118e46983188d80 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Tue, 8 Sep 2026 13:55:17 -0700
Subject: [PATCH 4/4] Test nullable global memory aggregations with dictionary
 and mixed values

No existing groupby test combined null keys with dictionary values, so
the nullable dictionary kernels and the per-type double launch for mixed
value tables never ran under test. Add tests that compare the hash path
against the sort-based implementation for dictionary-only,
non-dictionary-only, and mixed value tables, on both the sparse and
dense output paths, with null keys excluded and included.
---
 cpp/tests/CMakeLists.txt                      |   1 +
 .../groupby/global_memory_aggs_tests.cpp      | 164 ++++++++++++++++++
 2 files changed, 165 insertions(+)
 create mode 100644 cpp/tests/groupby/global_memory_aggs_tests.cpp

diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt
index 9153e1cbb82d..f79882c37d49 100644
--- a/cpp/tests/CMakeLists.txt
+++ b/cpp/tests/CMakeLists.txt
@@ -126,6 +126,7 @@ ConfigureTest(
   groupby/count_scan_tests.cpp
   groupby/count_tests.cpp
   groupby/covariance_tests.cpp
+  groupby/global_memory_aggs_tests.cpp
   groupby/groupby_test_util.cpp
   groupby/groups_tests.cpp
   groupby/histogram_tests.cpp
diff --git a/cpp/tests/groupby/global_memory_aggs_tests.cpp b/cpp/tests/groupby/global_memory_aggs_tests.cpp
new file mode 100644
index 000000000000..6de4b961a9db
--- /dev/null
+++ b/cpp/tests/groupby/global_memory_aggs_tests.cpp
@@ -0,0 +1,164 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <cudf_test/base_fixture.hpp>
+#include <cudf_test/column_utilities.hpp>
+#include <cudf_test/column_wrapper.hpp>
+#include <cudf_test/table_utilities.hpp>
+
+#include <cudf/aggregation.hpp>
+#include <cudf/column/column_view.hpp>
+#include <cudf/copying.hpp>
+#include <cudf/groupby.hpp>
+#include <cudf/sorting.hpp>
+#include <cudf/table/table.hpp>
+#include <cudf/table/table_view.hpp>
+#include <cudf/types.hpp>
+#include <cudf/utilities/error.hpp>
+
+#include <cstddef>
+#include <memory>
+#include <vector>
+
+/*
+ * The hash groupby uses its global memory kernels whenever a value column is a dictionary or a
+ * SUM_OVERFLOW aggregation is requested.  Null keys select the nullable hash set, and more than two
+ * single-pass aggregations select the dense output kernel instead of the sparse one.  The nullable
+ * kernels are split by value column type, so a table with both dictionary and non-dictionary
+ * values launches once per type.  These tests run every combination of those paths against the
+ * sort-based implementation, with null keys both excluded and included.
+ */
+
+namespace {
+
+struct request_spec {
+  cudf::column_view values;
+  std::vector<cudf::aggregation::Kind> kinds;
+};
+
+std::unique_ptr<cudf::groupby_aggregation> make_aggregation(cudf::aggregation::Kind kind)
+{
+  switch (kind) {
+    case cudf::aggregation::SUM: return cudf::make_sum_aggregation<cudf::groupby_aggregation>();
+    case cudf::aggregation::MIN: return cudf::make_min_aggregation<cudf::groupby_aggregation>();
+    case cudf::aggregation::MAX: return cudf::make_max_aggregation<cudf::groupby_aggregation>();
+    case cudf::aggregation::COUNT_VALID:
+      return cudf::make_count_aggregation<cudf::groupby_aggregation>();
+    case cudf::aggregation::SUM_OVERFLOW:
+      return cudf::make_sum_overflow_aggregation<cudf::groupby_aggregation>();
+    default: CUDF_FAIL("Aggregation kind is not used by these tests.");
+  }
+}
+
+enum class implementation { HASH, SORT };
+
+// Result rows sorted by key, so the unordered hash output can be compared with the sorted output.
+struct sorted_result {
+  std::unique_ptr<cudf::table> keys;
+  std::vector<std::unique_ptr<cudf::table>> values;
+};
+
+sorted_result aggregate(cudf::column_view const& keys,
+                        std::vector<request_spec> const& specs,
+                        cudf::null_policy null_handling,
+                        implementation impl)
+{
+  std::vector<cudf::groupby::aggregation_request> requests;
+  for (auto const& spec : specs) {
+    auto& request  = requests.emplace_back();
+    request.values = spec.values;
+    for (auto kind : spec.kinds) {
+      request.aggregations.push_back(make_aggregation(kind));
+    }
+  }
+  // nth_element is not a hash aggregation, so adding it forces the sort-based implementation.  It
+  // is appended last so the positions of the requested results are unchanged.
+  if (impl == implementation::SORT) {
+    requests.front().aggregations.push_back(
+      cudf::make_nth_element_aggregation<cudf::groupby_aggregation>(0));
+  }
+
+  cudf::groupby::groupby gb_obj(cudf::table_view{{keys}}, null_handling);
+  auto [result_keys, results] = gb_obj.aggregate(requests);
+
+  auto const order = cudf::sorted_order(result_keys->view());
+  sorted_result sorted;
+  sorted.keys = cudf::gather(result_keys->view(), *order);
+  for (std::size_t i = 0; i < specs.size(); ++i) {
+    for (std::size_t j = 0; j < specs[i].kinds.size(); ++j) {
+      sorted.values.push_back(
+        cudf::gather(cudf::table_view{{results[i].results[j]->view()}}, *order));
+    }
+  }
+  return sorted;
+}
+
+void expect_hash_matches_sort(cudf::column_view const& keys, std::vector<request_spec> const& specs)
+{
+  for (auto null_handling : {cudf::null_policy::EXCLUDE, cudf::null_policy::INCLUDE}) {
+    auto const hash = aggregate(keys, specs, null_handling, implementation::HASH);
+    auto const sort = aggregate(keys, specs, null_handling, implementation::SORT);
+    CUDF_TEST_EXPECT_TABLES_EQUIVALENT(sort.keys->view(), hash.keys->view());
+    ASSERT_EQ(sort.values.size(), hash.values.size());
+    for (std::size_t i = 0; i < sort.values.size(); ++i) {
+      CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(sort.values[i]->get_column(0),
+                                          hash.values[i]->get_column(0));
+    }
+  }
+}
+
+}  // namespace
+
+struct groupby_global_memory_aggs_test : public cudf::test::BaseFixture {
+  // Two null keys and one null in each value column, so null handling is exercised on both sides
+  // of every aggregation.
+  cudf::test::fixed_width_column_wrapper<int32_t> keys{{1, 2, 3, 1, 2, 2, 1, 3, 3, 2},
+                                                       {1, 0, 1, 1, 1, 1, 0, 1, 1, 1}};
+  cudf::test::dictionary_column_wrapper<int32_t> dictionary_values{
+    {10, 20, 30, 40, 50, 60, 70, 80, 90, 100}, {1, 1, 1, 1, 0, 1, 1, 1, 1, 1}};
+  cudf::test::fixed_width_column_wrapper<int64_t> plain_values{{5, 4, 3, 2, 1, 0, -1, -2, -3, -4},
+                                                               {1, 1, 1, 1, 1, 1, 1, 1, 0, 1}};
+};
+
+TEST_F(groupby_global_memory_aggs_test, DictionaryOnlySparse)
+{
+  expect_hash_matches_sort(keys, {{dictionary_values, {cudf::aggregation::SUM}}});
+}
+
+TEST_F(groupby_global_memory_aggs_test, DictionaryOnlyDense)
+{
+  expect_hash_matches_sort(
+    keys,
+    {{dictionary_values,
+      {cudf::aggregation::SUM, cudf::aggregation::MIN, cudf::aggregation::MAX}}});
+}
+
+TEST_F(groupby_global_memory_aggs_test, NonDictionaryOnlySparse)
+{
+  expect_hash_matches_sort(keys, {{plain_values, {cudf::aggregation::SUM_OVERFLOW}}});
+}
+
+TEST_F(groupby_global_memory_aggs_test, NonDictionaryOnlyDense)
+{
+  expect_hash_matches_sort(
+    keys,
+    {{plain_values,
+      {cudf::aggregation::SUM_OVERFLOW, cudf::aggregation::MIN, cudf::aggregation::MAX}}});
+}
+
+TEST_F(groupby_global_memory_aggs_test, MixedSparse)
+{
+  expect_hash_matches_sort(
+    keys,
+    {{dictionary_values, {cudf::aggregation::SUM}}, {plain_values, {cudf::aggregation::MAX}}});
+}
+
+TEST_F(groupby_global_memory_aggs_test, MixedDense)
+{
+  expect_hash_matches_sort(
+    keys,
+    {{dictionary_values, {cudf::aggregation::SUM, cudf::aggregation::MIN}},
+     {plain_values, {cudf::aggregation::MAX, cudf::aggregation::COUNT_VALID}}});
+}