Reduce compute_global_memory_aggs_null build time

cpp/CMakeLists.txt

@@ -667,6 +667,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

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)));
   }
 };
 

cpp/src/groupby/hash/compute_global_memory_aggs.hpp

@@ -5,10 +5,12 @@
 #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/device_uvector.hpp>
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/stream>

cpp/src/groupby/hash/compute_global_memory_aggs_null.cu

@@ -3,20 +3,45 @@
  * 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,
                                                   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);
+                                                  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

cpp/src/groupby/hash/compute_global_memory_aggs_null.hpp

@@ -0,0 +1,56 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * 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 {
+
+/*
+ * 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,
+                                      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_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,
+                                       cuda::stream_ref stream,
+                                       rmm::device_async_resource_ref mr);
+
+}  // namespace cudf::groupby::detail::hash

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,
+                                      cuda::stream_ref 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

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,
+                                  cuda::stream_ref stream)
+{
+  launch_null_dense_filtered<true>(
+    target_indices, aggs, input_values, output_values, num_items, stream);
+}
+
+}  // namespace cudf::groupby::detail::hash

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,
+                                      cuda::stream_ref stream)
+{
+  launch_null_dense_filtered<false>(
+    target_indices, aggs, input_values, output_values, num_items, stream);
+}
+
+}  // namespace cudf::groupby::detail::hash

cpp/src/groupby/hash/compute_global_memory_aggs_null_kernels.cuh

@@ -0,0 +1,74 @@
+/*
+ * 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)
+{
+  // 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>
+void launch_filtered_single_pass_aggs(Index num_items, Function fn, cuda::stream_ref 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.get()>>>(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,
+                                 cuda::stream_ref 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,
+                                cuda::stream_ref 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