Fix radix_sort tuning namespace (#6755)

Author: bernhardmgruber

c/parallel/src/radix_sort.cu

@@ -306,7 +306,7 @@ CUresult cccl_device_radix_sort_build_ex(
     check(cccl_type_name_from_nvrtc<OffsetT>(&offset_t));
 
     const auto policy_hub_expr =
-      std::format("cub::detail::radix::policy_hub<{}, {}, {}>", key_cpp, value_cpp, offset_t);
+      std::format("cub::detail::radix_sort::policy_hub<{}, {}, {}>", key_cpp, value_cpp, offset_t);
 
     const std::string final_src = std::format(
       R"XXX(
@@ -326,7 +326,7 @@ using {5} = {6}::MaxPolicy;
 #include <cub/detail/ptx-json/json.cuh>
 __device__ consteval auto& policy_generator() {{
   return ptx_json::id<ptx_json::string("device_radix_sort_policy")>()
-    = cub::detail::radix::RadixSortPolicyWrapper<{5}::ActivePolicy>::EncodedPolicy();
+    = cub::detail::radix_sort::RadixSortPolicyWrapper<{5}::ActivePolicy>::EncodedPolicy();
 }}
 )XXX",
       input_keys_it.value_type.size, // 0

cub/cub/device/dispatch/dispatch_radix_sort.cuh

@@ -115,7 +115,7 @@ template <SortOrder Order,
           typename ValueT,
           typename OffsetT,
           typename DecomposerT  = detail::identity_decomposer_t,
-          typename PolicyHub    = detail::radix::policy_hub<KeyT, ValueT, OffsetT>,
+          typename PolicyHub    = detail::radix_sort::policy_hub<KeyT, ValueT, OffsetT>,
           typename KernelSource = detail::radix_sort::
             DeviceRadixSortKernelSource<typename PolicyHub::MaxPolicy, Order, KeyT, ValueT, OffsetT, DecomposerT>,
           typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY>
@@ -925,7 +925,7 @@ struct DispatchRadixSort
   template <typename ActivePolicyT>
   CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t Invoke(ActivePolicyT policy = {})
   {
-    auto wrapped_policy = detail::radix::MakeRadixSortPolicyWrapper(policy);
+    auto wrapped_policy = detail::radix_sort::MakeRadixSortPolicyWrapper(policy);
 
     // Return if empty problem, or if no bits to sort and double-buffering is used
     if (num_items == 0 || (begin_bit == end_bit && is_overwrite_okay))

cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh

@@ -124,7 +124,7 @@ template <SortOrder Order,
           typename BeginOffsetIteratorT,
           typename EndOffsetIteratorT,
           typename SegmentSizeT,
-          typename PolicyHub    = detail::radix::policy_hub<KeyT, ValueT, SegmentSizeT>,
+          typename PolicyHub    = detail::radix_sort::policy_hub<KeyT, ValueT, SegmentSizeT>,
           typename DecomposerT  = detail::identity_decomposer_t,
           typename KernelSource = detail::radix_sort::DeviceSegmentedRadixSortKernelSource<
             typename PolicyHub::MaxPolicy,
@@ -517,7 +517,7 @@ struct DispatchSegmentedRadixSort
     // Force kernel code-generation in all compiler passes
     return InvokePasses(kernel_source.SegmentedRadixSortKernel(),
                         kernel_source.AltSegmentedRadixSortKernel(),
-                        detail::radix::MakeRadixSortPolicyWrapper(policy));
+                        detail::radix_sort::MakeRadixSortPolicyWrapper(policy));
   }
 
   //------------------------------------------------------------------------------

cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh

@@ -22,7 +22,7 @@
 
 CUB_NAMESPACE_BEGIN
 
-namespace detail::radix
+namespace detail::radix_sort
 {
 // sm90 default
 template <size_t KeySize, size_t ValueSize, size_t OffsetSize>
@@ -1040,6 +1040,6 @@ struct policy_hub
 
   using MaxPolicy = Policy1000;
 };
-} // namespace detail::radix
+} // namespace detail::radix_sort
 
 CUB_NAMESPACE_END