diff --git a/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh b/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh
index a94a19b17b0..8609862c2cd 100644
--- a/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh
@@ -101,7 +101,7 @@ struct policy_hub
 {
   using ValueT = it_value_t<InputIteratorT>;
 
-  struct Policy500 : ChainedPolicy<500, Policy500, Policy500>
+  struct Policy500 : detail::chained_policy<500, Policy500, Policy500>
   {
     using AdjacentDifferencePolicy =
       agent_adjacent_difference_policy<128,
diff --git a/cub/cub/device/dispatch/tuning/tuning_for.cuh b/cub/cub/device/dispatch/tuning/tuning_for.cuh
index e0537dd4f40..ebcf7549a24 100644
--- a/cub/cub/device/dispatch/tuning/tuning_for.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_for.cuh
@@ -66,7 +66,7 @@ struct policy_selector
 // TODO(bgruber): remove once we publish the tuning API
 struct policy_hub_t
 {
-  struct policy_500_t : ChainedPolicy<500, policy_500_t, policy_500_t>
+  struct policy_500_t : detail::chained_policy<500, policy_500_t, policy_500_t>
   {
     using for_policy_t = policy_t<256, 2>;
   };
diff --git a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh
index 20bdf1578b1..41251ff4339 100644
--- a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh
@@ -214,7 +214,7 @@ struct policy_hub
   }
 
   // SM50
-  struct Policy500 : ChainedPolicy<500, Policy500, Policy500>
+  struct Policy500 : detail::chained_policy<500, Policy500, Policy500>
   {
     // TODO This might be worth it to separate usual histogram and the multi one
     using AgentHistogramPolicyT =
@@ -222,7 +222,7 @@ struct policy_hub
   };
 
   // SM90
-  struct Policy900 : ChainedPolicy<900, Policy900, Policy500>
+  struct Policy900 : detail::chained_policy<900, Policy900, Policy500>
   {
     // Use values from tuning if a specialization exists, otherwise pick Policy500
     template <typename Tuning>
@@ -245,7 +245,7 @@ struct policy_hub
     static constexpr int init_kernel_pdl_trigger_max_bins = 2048;
   };
 
-  struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy900>
+  struct Policy1000 : detail::chained_policy<1000, Policy1000, Policy900>
   {
     // Use values from tuning if a specialization exists, otherwise pick Policy900
     template <typename Tuning>
diff --git a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh
index 94928b0de4f..e5e6120eade 100644
--- a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh
@@ -102,7 +102,7 @@ struct policy_hub
 {
   using KeyT = it_value_t<KeyIteratorT>;
 
-  struct Policy500 : ChainedPolicy<500, Policy500, Policy500>
+  struct Policy500 : detail::chained_policy<500, Policy500, Policy500>
   {
     using MergeSortPolicy =
       agent_merge_sort_policy<256,
@@ -116,7 +116,7 @@ struct policy_hub
 #if defined(_NVHPC_CUDA)
   using Policy520 = Policy500;
 #else
-  struct Policy520 : ChainedPolicy<520, Policy520, Policy500>
+  struct Policy520 : detail::chained_policy<520, Policy520, Policy500>
   {
     using MergeSortPolicy =
       agent_merge_sort_policy<512,
@@ -127,7 +127,7 @@ struct policy_hub
   };
 #endif
 
-  struct Policy600 : ChainedPolicy<600, Policy600, Policy520>
+  struct Policy600 : detail::chained_policy<600, Policy600, Policy520>
   {
     using MergeSortPolicy =
       agent_merge_sort_policy<256,
diff --git a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh
index dae9bd75e6e..ce3ce34ddf1 100644
--- a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh
@@ -1037,7 +1037,7 @@ struct policy_hub
   //------------------------------------------------------------------------------
 
   /// SM50
-  struct Policy500 : ChainedPolicy<500, Policy500, Policy500>
+  struct Policy500 : detail::chained_policy<500, Policy500, Policy500>
   {
     static constexpr int PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5; // 3.5B 32b keys/s, 1.92B 32b pairs/s (TitanX)
     static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5;
@@ -1129,7 +1129,7 @@ struct policy_hub
   };
 
   /// SM60 (GP100)
-  struct Policy600 : ChainedPolicy<600, Policy600, Policy500>
+  struct Policy600 : detail::chained_policy<600, Policy600, Policy500>
   {
     static constexpr int PRIMARY_RADIX_BITS     = (sizeof(KeyT) > 1) ? 7 : 5; // 6.9B 32b keys/s (Quadro P100)
     static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5;
@@ -1222,7 +1222,7 @@ struct policy_hub
   };
 
   /// SM61 (GP104)
-  struct Policy610 : ChainedPolicy<610, Policy610, Policy600>
+  struct Policy610 : detail::chained_policy<610, Policy610, Policy600>
   {
     static constexpr int PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5; // 3.4B 32b keys/s, 1.83B 32b pairs/s (1080)
     static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5;
@@ -1315,7 +1315,7 @@ struct policy_hub
   };
 
   /// SM62 (Tegra, less RF)
-  struct Policy620 : ChainedPolicy<620, Policy620, Policy610>
+  struct Policy620 : detail::chained_policy<620, Policy620, Policy610>
   {
     static constexpr int PRIMARY_RADIX_BITS  = 5;
     static constexpr int ALT_RADIX_BITS      = PRIMARY_RADIX_BITS - 1;
@@ -1390,7 +1390,7 @@ struct policy_hub
   };
 
   /// SM70 (GV100)
-  struct Policy700 : ChainedPolicy<700, Policy700, Policy620>
+  struct Policy700 : detail::chained_policy<700, Policy700, Policy620>
   {
     static constexpr int PRIMARY_RADIX_BITS     = (sizeof(KeyT) > 1) ? 7 : 5; // 7.62B 32b keys/s (GV100)
     static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5;
@@ -1484,7 +1484,7 @@ struct policy_hub
   };
 
   /// SM80
-  struct Policy800 : ChainedPolicy<800, Policy800, Policy700>
+  struct Policy800 : detail::chained_policy<800, Policy800, Policy700>
   {
     static constexpr int PRIMARY_RADIX_BITS     = (sizeof(KeyT) > 1) ? 7 : 5;
     static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5;
@@ -1699,12 +1699,12 @@ struct policy_hub
   };
 
   struct Policy900
-      : ChainedPolicy<900, Policy900, Policy800>
+      : detail::chained_policy<900, Policy900, Policy800>
       , OnesweepSmallKeyTunedPolicy<sm90_small_key_tuning<sizeof(KeyT), KEYS_ONLY ? 0 : sizeof(ValueT), sizeof(OffsetT)>>
   {};
 
   struct Policy1000
-      : ChainedPolicy<1000, Policy1000, Policy900>
+      : detail::chained_policy<1000, Policy1000, Policy900>
       , OnesweepSmallKeyTunedPolicy<
           sm100_small_key_tuning<ValueT, sizeof(KeyT), KEYS_ONLY ? 0 : sizeof(ValueT), sizeof(OffsetT)>>
   {};
diff --git a/cub/cub/device/dispatch/tuning/tuning_reduce.cuh b/cub/cub/device/dispatch/tuning/tuning_reduce.cuh
index c2f262943ae..a3d0df5ddaa 100644
--- a/cub/cub/device/dispatch/tuning/tuning_reduce.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_reduce.cuh
@@ -251,7 +251,7 @@ get_sm100_tuning(type_t accum_t, op_kind_t operation_t, int offset_size, int acc
 template <typename AccumT, typename OffsetT, typename ReductionOpT>
 struct policy_hub
 {
-  struct Policy500 : ChainedPolicy<500, Policy500, Policy500>
+  struct Policy500 : detail::chained_policy<500, Policy500, Policy500>
   {
     static constexpr int threads_per_block  = 256;
     static constexpr int items_per_thread   = 20;
@@ -279,7 +279,7 @@ struct policy_hub
                           NoScaling<ReducePolicy::BLOCK_THREADS, ReducePolicy::ITEMS_PER_THREAD>>;
   };
 
-  struct Policy600 : ChainedPolicy<600, Policy600, Policy500>
+  struct Policy600 : detail::chained_policy<600, Policy600, Policy500>
   {
     static constexpr int threads_per_block  = 256;
     static constexpr int items_per_thread   = 16;
@@ -307,7 +307,7 @@ struct policy_hub
                           NoScaling<ReducePolicy::BLOCK_THREADS, ReducePolicy::ITEMS_PER_THREAD>>;
   };
 
-  struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy600>
+  struct Policy1000 : detail::chained_policy<1000, Policy1000, Policy600>
   {
     // Use values from tuning if a specialization exists, otherwise pick Policy600
     template <typename Tuning>
diff --git a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh
index c614c67d7b6..097d35c6b8d 100644
--- a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh
@@ -920,7 +920,7 @@ struct policy_hub
 
   struct Policy500
       : DefaultPolicy500
-      , ChainedPolicy<500, Policy500, Policy500>
+      , detail::chained_policy<500, Policy500, Policy500>
   {};
 
   // Use values from tuning if a specialization exists, otherwise pick DefaultPolicy
@@ -936,7 +936,7 @@ struct policy_hub
   template <typename Tuning>
   static auto select_agent_policy(long) -> typename DefaultPolicy<LOAD_DEFAULT>::ReduceByKeyPolicyT;
 
-  struct Policy800 : ChainedPolicy<800, Policy800, Policy500>
+  struct Policy800 : detail::chained_policy<800, Policy800, Policy500>
   {
     using ReduceByKeyPolicyT =
       decltype(select_agent_policy<sm80_tuning<KeyT, AccumT, is_primitive_op<ReductionOpT>()>>(0));
@@ -947,16 +947,16 @@ struct policy_hub
 
   struct Policy860
       : DefaultPolicy860
-      , ChainedPolicy<860, Policy860, Policy800>
+      , detail::chained_policy<860, Policy860, Policy800>
   {};
 
-  struct Policy900 : ChainedPolicy<900, Policy900, Policy860>
+  struct Policy900 : detail::chained_policy<900, Policy900, Policy860>
   {
     using ReduceByKeyPolicyT =
       decltype(select_agent_policy<sm90_tuning<KeyT, AccumT, is_primitive_op<ReductionOpT>()>>(0));
   };
 
-  struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy900>
+  struct Policy1000 : detail::chained_policy<1000, Policy1000, Policy900>
   {
     // Use values from tuning if a specialization exists, otherwise fall back to SM90
     template <typename Tuning>
diff --git a/cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh b/cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh
index 30b3e343a95..cd8437184cd 100644
--- a/cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh
@@ -293,7 +293,7 @@ struct policy_hub
 
   struct Policy500
       : DefaultPolicy500
-      , ChainedPolicy<500, Policy500, Policy500>
+      , detail::chained_policy<500, Policy500, Policy500>
   {};
 
   // Use values from tuning if a specialization exists, otherwise pick the default
@@ -308,7 +308,7 @@ struct policy_hub
   template <typename Tuning>
   static auto select_agent_policy(long) -> typename DefaultPolicy<LOAD_DEFAULT>::ReduceByKeyPolicyT;
 
-  struct Policy800 : ChainedPolicy<800, Policy800, Policy500>
+  struct Policy800 : detail::chained_policy<800, Policy800, Policy500>
   {
     using ReduceByKeyPolicyT = decltype(select_agent_policy<sm80_tuning<LengthT, KeyT>>(0));
   };
@@ -318,15 +318,15 @@ struct policy_hub
 
   struct Policy860
       : DefaultPolicy860
-      , ChainedPolicy<860, Policy860, Policy800>
+      , detail::chained_policy<860, Policy860, Policy800>
   {};
 
-  struct Policy900 : ChainedPolicy<900, Policy900, Policy860>
+  struct Policy900 : detail::chained_policy<900, Policy900, Policy860>
   {
     using ReduceByKeyPolicyT = decltype(select_agent_policy<sm90_tuning<LengthT, KeyT>>(0));
   };
 
-  struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy900>
+  struct Policy1000 : detail::chained_policy<1000, Policy1000, Policy900>
   {
     // Use values from tuning if a specialization exists, otherwise pick Policy900
     template <typename Tuning>
diff --git a/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh b/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh
index f2329043af2..38745f93030 100644
--- a/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh
@@ -313,7 +313,7 @@ struct policy_hub
 
   struct Policy500
       : DefaultPolicy500
-      , ChainedPolicy<500, Policy500, Policy500>
+      , detail::chained_policy<500, Policy500, Policy500>
   {};
 
   // Use values from tuning if a specialization exists, otherwise pick the default
@@ -330,7 +330,7 @@ struct policy_hub
   static auto select_agent_policy(long) ->
     typename DefaultPolicy<BLOCK_LOAD_WARP_TRANSPOSE, LengthT, LOAD_DEFAULT>::RleSweepPolicyT;
 
-  struct Policy800 : ChainedPolicy<800, Policy800, Policy500>
+  struct Policy800 : detail::chained_policy<800, Policy800, Policy500>
   {
     using RleSweepPolicyT = decltype(select_agent_policy<sm80_tuning<LengthT, KeyT>>(0));
   };
@@ -341,15 +341,15 @@ struct policy_hub
 
   struct Policy860
       : DefaultPolicy860
-      , ChainedPolicy<860, Policy860, Policy800>
+      , detail::chained_policy<860, Policy860, Policy800>
   {};
 
-  struct Policy900 : ChainedPolicy<900, Policy900, Policy860>
+  struct Policy900 : detail::chained_policy<900, Policy900, Policy860>
   {
     using RleSweepPolicyT = decltype(select_agent_policy<sm90_tuning<LengthT, KeyT>>(0));
   };
 
-  struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy900>
+  struct Policy1000 : detail::chained_policy<1000, Policy1000, Policy900>
   {
     // Use values from tuning if a specialization exists, otherwise pick Policy900
     template <typename Tuning>
diff --git a/cub/cub/device/dispatch/tuning/tuning_scan.cuh b/cub/cub/device/dispatch/tuning/tuning_scan.cuh
index 3e9204a036e..22cf4b4b4e8 100644
--- a/cub/cub/device/dispatch/tuning/tuning_scan.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_scan.cuh
@@ -590,13 +590,13 @@ struct policy_hub
   static constexpr BlockStoreAlgorithm scan_transposed_store =
     large_values ? BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED : BLOCK_STORE_WARP_TRANSPOSE;
 
-  struct Policy500 : ChainedPolicy<500, Policy500, Policy500>
+  struct Policy500 : detail::chained_policy<500, Policy500, Policy500>
   {
     // GTX Titan: 29.5B items/s (232.4 GB/s) @ 48M 32-bit T
     using ScanPolicyT =
       agent_scan_policy<128, 12, AccumT, BLOCK_LOAD_DIRECT, LOAD_CA, BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED, BLOCK_SCAN_RAKING>;
   };
-  struct Policy520 : ChainedPolicy<520, Policy520, Policy500>
+  struct Policy520 : detail::chained_policy<520, Policy520, Policy500>
   {
     // Titan X: 32.47B items/s @ 48M 32-bit T
     using ScanPolicyT =
@@ -611,7 +611,7 @@ struct policy_hub
 
   struct Policy600
       : DefaultPolicy
-      , ChainedPolicy<600, Policy600, Policy520>
+      , detail::chained_policy<600, Policy600, Policy520>
   {};
 
   // Use values from tuning if a specialization exists, otherwise pick DefaultPolicy
@@ -629,7 +629,7 @@ struct policy_hub
   template <typename Tuning>
   _CCCL_HOST_DEVICE static auto select_agent_policy(long) -> typename DefaultPolicy::ScanPolicyT;
 
-  struct Policy750 : ChainedPolicy<750, Policy750, Policy600>
+  struct Policy750 : detail::chained_policy<750, Policy750, Policy600>
   {
     // Use values from tuning if a specialization exists that matches a benchmark, otherwise pick Policy600
     template <typename Tuning,
@@ -659,7 +659,7 @@ struct policy_hub
       decltype(select_agent_policy750<sm75_tuning<InputValueT, AccumT, OffsetT, classify_op<ScanOpT>>, InputValueT>(0));
   };
 
-  struct Policy800 : ChainedPolicy<800, Policy800, Policy750>
+  struct Policy800 : detail::chained_policy<800, Policy800, Policy750>
   {
     using ScanPolicyT =
       decltype(select_agent_policy<sm80_tuning<classify_type<AccumT>,
@@ -670,15 +670,15 @@ struct policy_hub
 
   struct Policy860
       : DefaultPolicy
-      , ChainedPolicy<860, Policy860, Policy800>
+      , detail::chained_policy<860, Policy860, Policy800>
   {};
 
-  struct Policy900 : ChainedPolicy<900, Policy900, Policy860>
+  struct Policy900 : detail::chained_policy<900, Policy900, Policy860>
   {
     using ScanPolicyT = decltype(select_agent_policy<sm90_tuning<AccumT, is_primitive_op<ScanOpT>()>>(0));
   };
 
-  struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy900>
+  struct Policy1000 : detail::chained_policy<1000, Policy1000, Policy900>
   {
     // Use values from tuning if a specialization exists that matches a benchmark, otherwise pick Policy900
     template <typename Tuning,
diff --git a/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh
index 22fa39ab449..70efb3d70e7 100644
--- a/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh
@@ -935,7 +935,7 @@ struct policy_hub
   static constexpr int max_input_bytes      = static_cast<int>((::cuda::std::max) (sizeof(key_t), sizeof(AccumT)));
   static constexpr int combined_input_bytes = static_cast<int>(sizeof(key_t) + sizeof(AccumT));
 
-  struct Policy500 : ChainedPolicy<500, Policy500, Policy500>
+  struct Policy500 : detail::chained_policy<500, Policy500, Policy500>
   {
     static constexpr int nominal_4b_items_per_thread = 6;
     static constexpr int items_per_thread =
@@ -973,7 +973,7 @@ struct policy_hub
 
   struct Policy520
       : DefaultPolicy520
-      , ChainedPolicy<520, Policy520, Policy500>
+      , detail::chained_policy<520, Policy520, Policy500>
   {};
 
   // Use values from tuning if a specialization exists, otherwise pick the default
@@ -991,7 +991,7 @@ struct policy_hub
   // FIXME(bgruber): should we rather use `AccumT` instead of `ValueT` like the other default policies?
   static auto select_agent_policy(long) -> typename DefaultPolicy<LOAD_DEFAULT, ValueT>::ScanByKeyPolicyT;
 
-  struct Policy800 : ChainedPolicy<800, Policy800, Policy520>
+  struct Policy800 : detail::chained_policy<800, Policy800, Policy520>
   {
     using ScanByKeyPolicyT = decltype(select_agent_policy<sm80_tuning<key_t, ValueT, is_primitive_op<ScanOpT>()>>(0));
   };
@@ -1001,15 +1001,15 @@ struct policy_hub
 
   struct Policy860
       : DefaultPolicy860
-      , ChainedPolicy<860, Policy860, Policy800>
+      , detail::chained_policy<860, Policy860, Policy800>
   {};
 
-  struct Policy900 : ChainedPolicy<900, Policy900, Policy860>
+  struct Policy900 : detail::chained_policy<900, Policy900, Policy860>
   {
     using ScanByKeyPolicyT = decltype(select_agent_policy<sm90_tuning<key_t, ValueT, is_primitive_op<ScanOpT>()>>(0));
   };
 
-  struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy900>
+  struct Policy1000 : detail::chained_policy<1000, Policy1000, Policy900>
   {
     // Use values from tuning if a specialization exists, otherwise pick Policy900
     template <typename Tuning>
diff --git a/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh b/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh
index a1cf9919c59..d7ac127c86f 100644
--- a/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh
@@ -142,7 +142,7 @@ struct policy_selector_from_types
 template <typename AccumT, typename OffsetT, typename ReductionOpT>
 struct policy_hub
 {
-  struct Policy500 : ChainedPolicy<500, Policy500, Policy500>
+  struct Policy500 : detail::chained_policy<500, Policy500, Policy500>
   {
   private:
     static constexpr int items_per_vec_load = 4;
diff --git a/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh
index 36a504984fa..8fc42ce299c 100644
--- a/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh
@@ -405,7 +405,7 @@ struct policy_hub
   using DominantT                = ::cuda::std::_If<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>;
   static constexpr int KEYS_ONLY = ::cuda::std::is_same_v<ValueT, NullType>;
 
-  struct Policy500 : ChainedPolicy<500, Policy500, Policy500>
+  struct Policy500 : detail::chained_policy<500, Policy500, Policy500>
   {
     static constexpr int BLOCK_THREADS          = 256;
     static constexpr int RADIX_BITS             = sizeof(KeyT) > 1 ? 6 : 4;
@@ -438,7 +438,7 @@ struct policy_hub
                                        LOAD_DEFAULT>;
   };
 
-  struct Policy600 : ChainedPolicy<600, Policy600, Policy500>
+  struct Policy600 : detail::chained_policy<600, Policy600, Policy500>
   {
     static constexpr int BLOCK_THREADS          = 256;
     static constexpr int RADIX_BITS             = sizeof(KeyT) > 1 ? 6 : 4;
@@ -471,7 +471,7 @@ struct policy_hub
                                        LOAD_DEFAULT>;
   };
 
-  struct Policy610 : ChainedPolicy<610, Policy610, Policy600>
+  struct Policy610 : detail::chained_policy<610, Policy610, Policy600>
   {
     static constexpr int BLOCK_THREADS          = 256;
     static constexpr int RADIX_BITS             = sizeof(KeyT) > 1 ? 6 : 4;
@@ -504,7 +504,7 @@ struct policy_hub
                                        LOAD_DEFAULT>;
   };
 
-  struct Policy620 : ChainedPolicy<620, Policy620, Policy610>
+  struct Policy620 : detail::chained_policy<620, Policy620, Policy610>
   {
     static constexpr int BLOCK_THREADS          = 256;
     static constexpr int RADIX_BITS             = sizeof(KeyT) > 1 ? 5 : 4;
@@ -537,7 +537,7 @@ struct policy_hub
                                        LOAD_DEFAULT>;
   };
 
-  struct Policy700 : ChainedPolicy<700, Policy700, Policy620>
+  struct Policy700 : detail::chained_policy<700, Policy700, Policy620>
   {
     static constexpr int BLOCK_THREADS          = 256;
     static constexpr int RADIX_BITS             = sizeof(KeyT) > 1 ? 6 : 4;
@@ -570,7 +570,7 @@ struct policy_hub
                                        LOAD_DEFAULT>;
   };
 
-  struct Policy800 : ChainedPolicy<800, Policy800, Policy700>
+  struct Policy800 : detail::chained_policy<800, Policy800, Policy700>
   {
     static constexpr int BLOCK_THREADS          = 256;
     static constexpr int PARTITIONING_THRESHOLD = 500;
@@ -601,7 +601,7 @@ struct policy_hub
                                        LOAD_DEFAULT>;
   };
 
-  struct Policy860 : ChainedPolicy<860, Policy860, Policy800>
+  struct Policy860 : detail::chained_policy<860, Policy860, Policy800>
   {
     static constexpr int BLOCK_THREADS          = 256;
     static constexpr int PARTITIONING_THRESHOLD = 500;
diff --git a/cub/cub/device/dispatch/tuning/tuning_select_if.cuh b/cub/cub/device/dispatch/tuning/tuning_select_if.cuh
index b80d4e6d684..0f52b214ffa 100644
--- a/cub/cub/device/dispatch/tuning/tuning_select_if.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_select_if.cuh
@@ -1590,7 +1590,7 @@ struct policy_hub
 
   struct Policy500
       : DefaultPolicy500
-      , ChainedPolicy<500, Policy500, Policy500>
+      , detail::chained_policy<500, Policy500, Policy500>
   {};
 
   // Use values from tuning if a specialization exists, otherwise pick the default
@@ -1605,7 +1605,7 @@ struct policy_hub
   template <typename Tuning>
   static auto select_agent_policy(long) -> typename DefaultPolicy<LOAD_DEFAULT>::SelectIfPolicyT;
 
-  struct Policy800 : ChainedPolicy<800, Policy800, Policy500>
+  struct Policy800 : detail::chained_policy<800, Policy800, Policy500>
   {
     using SelectIfPolicyT =
       decltype(select_agent_policy<sm80_tuning<InputT,
@@ -1621,10 +1621,10 @@ struct policy_hub
 
   struct Policy860
       : DefaultPolicy860
-      , ChainedPolicy<860, Policy860, Policy800>
+      , detail::chained_policy<860, Policy860, Policy800>
   {};
 
-  struct Policy900 : ChainedPolicy<900, Policy900, Policy860>
+  struct Policy900 : detail::chained_policy<900, Policy900, Policy860>
   {
     using SelectIfPolicyT =
       decltype(select_agent_policy<sm90_tuning<InputT,
@@ -1635,7 +1635,7 @@ struct policy_hub
                                                classify_input_size<InputT>()>>(0));
   };
 
-  struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy900>
+  struct Policy1000 : detail::chained_policy<1000, Policy1000, Policy900>
   {
     // Use values from tuning if a specialization exists, otherwise pick Policy900
     template <typename Tuning>
diff --git a/cub/cub/device/dispatch/tuning/tuning_three_way_partition.cuh b/cub/cub/device/dispatch/tuning/tuning_three_way_partition.cuh
index f97cb8dbfef..15498a0b7f1 100644
--- a/cub/cub/device/dispatch/tuning/tuning_three_way_partition.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_three_way_partition.cuh
@@ -396,7 +396,7 @@ struct policy_hub
 
   struct Policy500
       : DefaultPolicy500
-      , ChainedPolicy<500, Policy500, Policy500>
+      , detail::chained_policy<500, Policy500, Policy500>
   {};
 
   // Use values from tuning if a specialization exists, otherwise pick DefaultPolicy
@@ -413,7 +413,7 @@ struct policy_hub
   static _CCCL_HOST_DEVICE auto select_agent_policy(long) -> typename DefaultPolicy<
     default_delay_constructor_t<typename accumulator_pack_t<OffsetT>::pack_t>>::ThreeWayPartitionPolicy;
 
-  struct Policy800 : ChainedPolicy<800, Policy800, Policy500>
+  struct Policy800 : detail::chained_policy<800, Policy800, Policy500>
   {
     using ThreeWayPartitionPolicy = decltype(select_agent_policy<sm80_tuning<InputT, OffsetT>>(0));
   };
@@ -423,15 +423,15 @@ struct policy_hub
 
   struct Policy860
       : DefaultPolicy860
-      , ChainedPolicy<860, Policy860, Policy800>
+      , detail::chained_policy<860, Policy860, Policy800>
   {};
 
-  struct Policy900 : ChainedPolicy<900, Policy900, Policy860>
+  struct Policy900 : detail::chained_policy<900, Policy900, Policy860>
   {
     using ThreeWayPartitionPolicy = decltype(select_agent_policy<sm90_tuning<InputT, OffsetT>>(0));
   };
 
-  struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy900>
+  struct Policy1000 : detail::chained_policy<1000, Policy1000, Policy900>
   {
     // Use values from tuning if a specialization exists, otherwise pick Policy900
     template <typename Tuning>
diff --git a/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh
index 8cb2c75f2f7..d43bba6d6e7 100644
--- a/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh
@@ -804,7 +804,7 @@ struct policy_hub
 
   struct Policy500
       : DefaultPolicy500
-      , ChainedPolicy<500, Policy500, Policy500>
+      , detail::chained_policy<500, Policy500, Policy500>
   {};
 
   // Use values from tuning if a specialization exists, otherwise pick the default
@@ -824,10 +824,10 @@ struct policy_hub
 
   struct Policy520
       : DefaultPolicy520
-      , ChainedPolicy<520, Policy520, Policy500>
+      , detail::chained_policy<520, Policy520, Policy500>
   {};
 
-  struct Policy800 : ChainedPolicy<800, Policy800, Policy520>
+  struct Policy800 : detail::chained_policy<800, Policy800, Policy520>
   {
     using UniqueByKeyPolicyT = decltype(select_agent_policy<sm80_tuning<KeyT, ValueT>>(0));
   };
@@ -837,15 +837,15 @@ struct policy_hub
 
   struct Policy860
       : DefaultPolicy860
-      , ChainedPolicy<860, Policy860, Policy800>
+      , detail::chained_policy<860, Policy860, Policy800>
   {};
 
-  struct Policy900 : ChainedPolicy<900, Policy900, Policy860>
+  struct Policy900 : detail::chained_policy<900, Policy900, Policy860>
   {
     using UniqueByKeyPolicyT = decltype(select_agent_policy<sm90_tuning<KeyT, ValueT>>(0));
   };
 
-  struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy900>
+  struct Policy1000 : detail::chained_policy<1000, Policy1000, Policy900>
   {
     // Use values from tuning if a specialization exists, otherwise pick Policy900
     template <typename Tuning>
diff --git a/cub/cub/util_device.cuh b/cub/cub/util_device.cuh
index 09caa064b2a..b21d6465a5b 100644
--- a/cub/cub/util_device.cuh
+++ b/cub/cub/util_device.cuh
@@ -628,10 +628,7 @@ _CCCL_HOST_DEVICE constexpr int LoadToSharedBufferSizeBytes(::cuda::std::size_t
   const int extra_space = (num_bytes == 0) ? 0 : detail::bulk_copy_min_align;
   return ::cuda::round_up(num_bytes, detail::bulk_copy_min_align) + extra_space;
 }
-} // namespace detail
 
-namespace detail
-{
 #if defined(CUB_DEFINE_RUNTIME_POLICIES)
 // TODO(bgruber): drop in CCCL 4.0 when we drop the dispatchers
 #  if !_CCCL_HAS_CONCEPTS()
@@ -689,15 +686,9 @@ __host__ __device__ constexpr PolicyT MakePolicyWrapper(PolicyT policy)
 {
   return policy;
 }
-} // namespace detail
-
-//----------------------------------------------------------------------------------------------------------------------
-// ChainedPolicy
 
 #if !_CCCL_COMPILER(NVRTC)
 
-namespace detail
-{
 // Forward declaration of the default kernel launcher factory
 struct TripleChevronFactory;
 
@@ -746,21 +737,18 @@ struct KernelConfig
     return launcher_factory.MaxSmOccupancy(sm_occupancy, kernel_ptr, threads_per_block);
   }
 };
-} // namespace detail
+
 #endif // !_CCCL_COMPILER(NVRTC)
 
-namespace detail
-{
 template <typename T>
 struct get_active_policy
 {
   using type = typename T::ActivePolicy;
 };
-} // namespace detail
 
 /// Helper for dispatching into a policy chain
 template <int PolicyPtxVersion, typename PolicyT, typename PrevPolicyT>
-struct ChainedPolicy
+struct chained_policy
 {
 private:
   static constexpr bool have_previous_policy = !::cuda::std::is_same_v<PolicyT, PrevPolicyT>;
@@ -799,7 +787,7 @@ public:
 
 private:
   template <int, typename, typename>
-  friend struct ChainedPolicy; // let us call find_and_invoke_policy of other ChainedPolicy instantiations
+  friend struct chained_policy; // let us call find_and_invoke_policy of other ChainedPolicy instantiations
 
 #if !_CCCL_COMPILER(NVRTC)
   template <int CcMult, int... CudaCcs, typename FunctorT>
@@ -832,6 +820,14 @@ private:
   }
 #endif // !_CCCL_COMPILER(NVRTC)
 };
+} // namespace detail
+
+/// Helper for dispatching into a policy chain
+/// Deprecated [Since 3.5]
+template <int PolicyPtxVersion, typename PolicyT, typename PrevPolicyT>
+using ChainedPolicy
+  CCCL_DEPRECATED_BECAUSE("Pass policy selectors into the environments of device-scope CUB algorithms to providing "
+                          "custom tunings.") = detail::chained_policy<PolicyPtxVersion, PolicyT, PrevPolicyT>;
 
 namespace detail
 {
diff --git a/cub/test/catch2_test_device_adjacent_difference_custom_policy_hub.cu b/cub/test/catch2_test_device_adjacent_difference_custom_policy_hub.cu
index 8eb222d5393..382ab9ed443 100644
--- a/cub/test/catch2_test_device_adjacent_difference_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_adjacent_difference_custom_policy_hub.cu
@@ -25,7 +25,7 @@ struct my_policy_hub
   using ValueT = cub::detail::it_value_t<InputIteratorT>;
 
   // from Policy500 of the CUB adjacent difference tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     using AdjacentDifferencePolicy =
       AgentAdjacentDifferencePolicy<128,
diff --git a/cub/test/catch2_test_device_histogram_custom_policy_hub.cu b/cub/test/catch2_test_device_histogram_custom_policy_hub.cu
index dd54ebaa1e9..d5cda9277af 100644
--- a/cub/test/catch2_test_device_histogram_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_histogram_custom_policy_hub.cu
@@ -21,7 +21,7 @@ template <class SampleT, class CounterT, int NumChannels, int NumActiveChannels,
 struct my_policy_hub
 {
   // simplified from Policy500 of the CUB histogram tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     using AgentHistogramPolicyT = AgentHistogramPolicy<384, 16, BLOCK_LOAD_DIRECT, LOAD_LDG, true, SMEM, false>;
     static constexpr int init_kernel_pdl_trigger_max_bins = 2048;
diff --git a/cub/test/catch2_test_device_merge_sort_custom_policy_hub.cu b/cub/test/catch2_test_device_merge_sort_custom_policy_hub.cu
index 5815c424434..5e18111b744 100644
--- a/cub/test/catch2_test_device_merge_sort_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_merge_sort_custom_policy_hub.cu
@@ -25,7 +25,7 @@ struct my_policy_hub
   using KeyT = cub::detail::it_value_t<KeyIteratorT>;
 
   // from Policy500 of the CUB merge sort tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     using MergeSortPolicy =
       AgentMergeSortPolicy<256,
diff --git a/cub/test/catch2_test_device_radix_sort_custom_policy_hub.cu b/cub/test/catch2_test_device_radix_sort_custom_policy_hub.cu
index 83bfe5fc2f1..1346a0e1a59 100644
--- a/cub/test/catch2_test_device_radix_sort_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_radix_sort_custom_policy_hub.cu
@@ -20,7 +20,7 @@ struct my_policy_hub
   using DominantT = KeyT;
 
   // from Policy500 of the CUB radix sort tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     static constexpr int PRIMARY_RADIX_BITS     = (sizeof(KeyT) > 1) ? 7 : 5;
     static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5;
diff --git a/cub/test/catch2_test_device_reduce_by_key_custom_policy_hub.cu b/cub/test/catch2_test_device_reduce_by_key_custom_policy_hub.cu
index 9537d0ea941..69bd7199c30 100644
--- a/cub/test/catch2_test_device_reduce_by_key_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_reduce_by_key_custom_policy_hub.cu
@@ -19,7 +19,7 @@ template <class ReductionOpT, class AccumT, class KeyT>
 struct my_policy_hub
 {
   // from Policy500 of the CUB reduce-by-key tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     using ReduceByKeyPolicyT =
       AgentReduceByKeyPolicy<128,
diff --git a/cub/test/catch2_test_device_reduce_dispatcher.cu b/cub/test/catch2_test_device_reduce_dispatcher.cu
index dd6d79a2f8f..fa31624edcb 100644
--- a/cub/test/catch2_test_device_reduce_dispatcher.cu
+++ b/cub/test/catch2_test_device_reduce_dispatcher.cu
@@ -21,7 +21,7 @@ using value_types = c2h::type_list<std::int8_t, std::int16_t, std::int32_t, std:
 template <typename AccumT>
 struct policy_hub_t
 {
-  struct policy_t : cub::ChainedPolicy<300, policy_t, policy_t>
+  struct policy_t : cub::detail::chained_policy<300, policy_t, policy_t>
   {
     static constexpr int threads_per_block  = 256;
     static constexpr int items_per_thread   = 16;
diff --git a/cub/test/catch2_test_device_rle_custom_policy_hub.cu b/cub/test/catch2_test_device_rle_custom_policy_hub.cu
index 1ed40a6a7d9..7ed9061490e 100644
--- a/cub/test/catch2_test_device_rle_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_rle_custom_policy_hub.cu
@@ -21,7 +21,7 @@ using namespace cub;
 template <typename LengthT, typename KeyT>
 struct my_policy_hub
 {
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     using RleSweepPolicyT =
       AgentRlePolicy<96,
diff --git a/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu b/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu
index 42a58c6a4c9..be53c0e935a 100644
--- a/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu
@@ -24,7 +24,7 @@ struct my_policy_hub
   using key_t = cub::detail::it_value_t<KeysInputIteratorT>;
 
   // from Policy500 of the CUB scan-by-key tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     using ScanByKeyPolicyT =
       AgentScanByKeyPolicy<128,
diff --git a/cub/test/catch2_test_device_scan_custom_policy_hub.cu b/cub/test/catch2_test_device_scan_custom_policy_hub.cu
index 813fc084923..913bb7a6e6d 100644
--- a/cub/test/catch2_test_device_scan_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_scan_custom_policy_hub.cu
@@ -23,7 +23,7 @@ template <typename InputValueT, typename OutputValueT, typename AccumT, typename
 struct my_policy_hub
 {
   // from Policy500 of the CUB scan tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     using ScanPolicyT =
       AgentScanPolicy<128, 12, AccumT, BLOCK_LOAD_DIRECT, LOAD_CA, BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED, BLOCK_SCAN_RAKING>;
diff --git a/cub/test/catch2_test_device_segmented_radix_sort_custom_policy_hub.cu b/cub/test/catch2_test_device_segmented_radix_sort_custom_policy_hub.cu
index 0bf62909ad5..92e84c3787f 100644
--- a/cub/test/catch2_test_device_segmented_radix_sort_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_segmented_radix_sort_custom_policy_hub.cu
@@ -21,7 +21,7 @@ template <typename KeyT, typename OffsetT>
 struct my_policy_hub
 {
   // from Policy500 of the CUB radix sort tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     static constexpr int PRIMARY_RADIX_BITS     = (sizeof(KeyT) > 1) ? 7 : 5;
     static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5;
diff --git a/cub/test/catch2_test_device_segmented_reduce_custom_policy_hub.cu b/cub/test/catch2_test_device_segmented_reduce_custom_policy_hub.cu
index 621d888eec6..0da58bc31e9 100644
--- a/cub/test/catch2_test_device_segmented_reduce_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_segmented_reduce_custom_policy_hub.cu
@@ -23,7 +23,7 @@ template <typename AccumT, typename OffsetT, typename ReductionOpT>
 struct my_policy_hub
 {
   // from Policy500 of the CUB segmented reduce tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     using ReducePolicy          = AgentReducePolicy<256, 20, AccumT, 4, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG>;
     using SingleTilePolicy      = ReducePolicy;
diff --git a/cub/test/catch2_test_device_segmented_sort_custom_policy_hub.cu b/cub/test/catch2_test_device_segmented_sort_custom_policy_hub.cu
index 3c2ba683509..4bf6be4ecbc 100644
--- a/cub/test/catch2_test_device_segmented_sort_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_segmented_sort_custom_policy_hub.cu
@@ -22,7 +22,7 @@ template <typename KeyT>
 struct my_policy_hub
 {
   // from Policy500 of the CUB segmented sort tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     static constexpr int BLOCK_THREADS          = 256;
     static constexpr int RADIX_BITS             = 6;
diff --git a/cub/test/catch2_test_device_select_if_custom_policy_hub.cu b/cub/test/catch2_test_device_select_if_custom_policy_hub.cu
index c88a5cbcf9d..1e024ef26e2 100644
--- a/cub/test/catch2_test_device_select_if_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_select_if_custom_policy_hub.cu
@@ -20,7 +20,7 @@ using namespace cub;
 template <class InputT>
 struct my_policy_hub
 {
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     static constexpr int nominal_4b_items_per_thread = 10;
     static constexpr int items_per_thread =
diff --git a/cub/test/catch2_test_device_select_unique_by_key_custom_policy_hub.cu b/cub/test/catch2_test_device_select_unique_by_key_custom_policy_hub.cu
index e4a524c9894..abd1e2c72f6 100644
--- a/cub/test/catch2_test_device_select_unique_by_key_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_select_unique_by_key_custom_policy_hub.cu
@@ -21,7 +21,7 @@ template <class KeyT, class ValueT>
 struct my_policy_hub
 {
   // from Policy500 of the CUB unique-by-key tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     using UniqueByKeyPolicyT =
       AgentUniqueByKeyPolicy<128,
diff --git a/cub/test/catch2_test_device_three_way_partition_custom_policy_hub.cu b/cub/test/catch2_test_device_three_way_partition_custom_policy_hub.cu
index 66e2a21d0de..4cd5e2f9580 100644
--- a/cub/test/catch2_test_device_three_way_partition_custom_policy_hub.cu
+++ b/cub/test/catch2_test_device_three_way_partition_custom_policy_hub.cu
@@ -22,7 +22,7 @@ template <class InputT, class OffsetT>
 struct my_policy_hub
 {
   // from Policy500 of the CUB three-way partition tunings
-  struct MaxPolicy : ChainedPolicy<500, MaxPolicy, MaxPolicy>
+  struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy>
   {
     using ThreeWayPartitionPolicy =
       AgentThreeWayPartitionPolicy<256,
diff --git a/cub/test/catch2_test_util_device.cu b/cub/test/catch2_test_util_device.cu
index 95316783bd2..e0d8975d2cc 100644
--- a/cub/test/catch2_test_util_device.cu
+++ b/cub/test/catch2_test_util_device.cu
@@ -94,10 +94,10 @@ C2H_TEST("PtxVersion returns a value from __CUDA_ARCH_LIST__/NV_TARGET_SM_INTEGE
 }
 #endif
 
-#define GEN_POLICY(cur, prev)                                             \
-  struct policy##cur : cub::ChainedPolicy<cur, policy##cur, policy##prev> \
-  {                                                                       \
-    static constexpr int value = cur;                                     \
+#define GEN_POLICY(cur, prev)                                                      \
+  struct policy##cur : cub::detail::chained_policy<cur, policy##cur, policy##prev> \
+  {                                                                                \
+    static constexpr int value = cur;                                              \
   }
 
 #ifdef CUDA_SM_LIST
@@ -174,6 +174,7 @@ check_chained_policy_prunes_to_cc_list(void* d_temp_storage, size_t& temp_storag
 
 DECLARE_LAUNCH_WRAPPER(check_chained_policy_prunes_to_cc_list, check_wrapper_all);
 
+// TODO(bgruber): drop in CCCL 4.0
 C2H_TEST("ChainedPolicy prunes based on __CUDA_ARCH_LIST__/NV_TARGET_SM_INTEGER_LIST", "[util][dispatch]")
 {
   check_wrapper_all();
@@ -263,6 +264,7 @@ struct policy_hub_minimal
   using max_policy = policy500;
 };
 
+// TODO(bgruber): drop in CCCL 4.0
 C2H_TEST("ChainedPolicy invokes correct policy", "[util][dispatch]")
 {
   SECTION("policy_hub_some")
diff --git a/cub/test/catch2_test_vsmem.cu b/cub/test/catch2_test_vsmem.cu
index 98b6c78aa83..2768205d31d 100644
--- a/cub/test/catch2_test_vsmem.cu
+++ b/cub/test/catch2_test_vsmem.cu
@@ -163,6 +163,8 @@ void __global__ __launch_bounds__(
   vsmem_helper_t::discard_temp_storage(temp_storage);
 }
 
+// TODO(bgruber): rewrite the below test to use policy selectors
+
 //----------------------------------------------------------------------------
 // Tuning policy chain
 //----------------------------------------------------------------------------
@@ -173,7 +175,7 @@ struct device_dummy_algorithm_policy_t
 
   static constexpr int FALLBACK_BLOCK_THREADS = 64;
 
-  struct policy_500 : cub::ChainedPolicy<500, policy_500, policy_500>
+  struct policy_500 : cub::detail::chained_policy<500, policy_500, policy_500>
   {
     using DummyAlgorithmPolicy = agent_dummy_algorithm_policy_t<256, cub::Nominal4BItemsToItems<item_t>(17)>;