From 95919f971e7cf13aacd650c324cc8c1468e67e12 Mon Sep 17 00:00:00 2001
From: "github-actions[bot]"
 <41898282+github-actions[bot]@users.noreply.github.com>
Date: Wed, 5 Feb 2025 16:22:17 +0100
Subject: [PATCH] Internalize cub::KernelConfig (#3683) (#3688)

(cherry picked from commit 32d05c747fa72a234dfe5816bbca55fbbd0f956a)

Co-authored-by: Federico Busato <50413820+fbusato@users.noreply.github.com>
---
 .../device/dispatch/dispatch_radix_sort.cuh   |  8 +++----
 cub/cub/device/dispatch/dispatch_reduce.cuh   |  6 ++++--
 cub/cub/util_device.cuh                       | 21 ++++++++-----------
 3 files changed, 17 insertions(+), 18 deletions(-)

diff --git a/cub/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_radix_sort.cuh
index d91392fef8a..6bd99bdf6c8 100644
--- a/cub/cub/device/dispatch/dispatch_radix_sort.cuh
+++ b/cub/cub/device/dispatch/dispatch_radix_sort.cuh
@@ -1184,11 +1184,11 @@ struct DispatchRadixSort
   struct PassConfig
   {
     UpsweepKernelT upsweep_kernel;
-    KernelConfig upsweep_config;
+    detail::KernelConfig upsweep_config;
     ScanKernelT scan_kernel;
-    KernelConfig scan_config;
+    detail::KernelConfig scan_config;
     DownsweepKernelT downsweep_kernel;
-    KernelConfig downsweep_config;
+    detail::KernelConfig downsweep_config;
     int radix_bits;
     int radix_digits;
     int max_downsweep_grid_size;
@@ -2135,7 +2135,7 @@ struct DispatchSegmentedRadixSort
   struct PassConfig
   {
     SegmentedKernelT segmented_kernel;
-    KernelConfig segmented_config;
+    detail::KernelConfig segmented_config;
     int radix_bits;
     int radix_digits;
 
diff --git a/cub/cub/device/dispatch/dispatch_reduce.cuh b/cub/cub/device/dispatch/dispatch_reduce.cuh
index aaa3c2ff95a..b39a1b5f4b3 100644
--- a/cub/cub/device/dispatch/dispatch_reduce.cuh
+++ b/cub/cub/device/dispatch/dispatch_reduce.cuh
@@ -474,7 +474,8 @@ struct DispatchReduce
       }
 
       // Init regular kernel configuration
-      KernelConfig reduce_config;
+      detail::KernelConfig reduce_config;
+      (void) reduce_config;
       error = CubDebug(reduce_config.Init(reduce_kernel, active_policy.Reduce(), launcher_factory));
       if (cudaSuccess != error)
       {
@@ -949,7 +950,8 @@ struct DispatchSegmentedReduce
       }
 
       // Init kernel configuration
-      KernelConfig segmented_reduce_config;
+      detail::KernelConfig segmented_reduce_config;
+      (void) &segmented_reduce_config;
       error =
         CubDebug(segmented_reduce_config.Init<typename ActivePolicyT::SegmentedReducePolicy>(segmented_reduce_kernel));
       if (cudaSuccess != error)
diff --git a/cub/cub/util_device.cuh b/cub/cub/util_device.cuh
index 76e5ccbaf01..3d98d7ace17 100644
--- a/cub/cub/util_device.cuh
+++ b/cub/cub/util_device.cuh
@@ -656,25 +656,18 @@ CUB_RUNTIME_FUNCTION detail::PolicyWrapper<PolicyT> MakePolicyWrapper(PolicyT po
 
 namespace detail
 {
+
 struct TripleChevronFactory;
-}
 
 /**
  * Kernel dispatch configuration
  */
 struct KernelConfig
 {
-  int block_threads;
-  int items_per_thread;
-  int tile_size;
-  int sm_occupancy;
-
-  CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE KernelConfig()
-      : block_threads(0)
-      , items_per_thread(0)
-      , tile_size(0)
-      , sm_occupancy(0)
-  {}
+  int block_threads{0};
+  int items_per_thread{0};
+  int tile_size{0};
+  int sm_occupancy{0};
 
   template <typename AgentPolicyT, typename KernelPtrT, typename LauncherFactory = detail::TripleChevronFactory>
   CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t
@@ -687,6 +680,10 @@ struct KernelConfig
   }
 };
 
+} // namespace detail
+
+using KernelConfig CCCL_DEPRECATED_BECAUSE("This class is considered an implementation detail") = detail::KernelConfig;
+
 /// Helper for dispatching into a policy chain
 template <int PolicyPtxVersion, typename PolicyT, typename PrevPolicyT>
 struct ChainedPolicy