diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh
index ab2abc62f15..edcac9c803a 100644
--- a/cub/cub/agent/agent_histogram.cuh
+++ b/cub/cub/agent/agent_histogram.cuh
@@ -43,7 +43,7 @@ enum BlockHistogramMemoryPreference
 #if _CCCL_HOSTED()
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(BlockHistogramMemoryPreference mempref) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(BlockHistogramMemoryPreference mempref) noexcept
 {
   switch (mempref)
   {
@@ -53,9 +53,8 @@ namespace detail
       return "SMEM";
     case BLEND:
       return "BLEND";
-    default:
-      return "<unknown BlockHistogramMemoryPreference>";
   }
+  return "<unknown BlockHistogramMemoryPreference>";
 }
 } // namespace detail
 
@@ -69,7 +68,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::BlockHistogramMemoryPreference, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::BlockHistogramMemoryPreference, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::BlockHistogramMemoryPreference& mempref, FmtCtx& ctx) const
diff --git a/cub/cub/agent/agent_radix_sort_onesweep.cuh b/cub/cub/agent/agent_radix_sort_onesweep.cuh
index 5980798d3db..ba994d90ce2 100644
--- a/cub/cub/agent/agent_radix_sort_onesweep.cuh
+++ b/cub/cub/agent/agent_radix_sort_onesweep.cuh
@@ -56,7 +56,7 @@ enum RadixSortStoreAlgorithm
 #if _CCCL_HOSTED()
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(RadixSortStoreAlgorithm algo) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(RadixSortStoreAlgorithm algo) noexcept
 {
   switch (algo)
   {
@@ -64,9 +64,8 @@ namespace detail
       return "RADIX_SORT_STORE_DIRECT";
     case RADIX_SORT_STORE_ALIGNED:
       return "RADIX_SORT_STORE_ALIGNED";
-    default:
-      return "<unknown RadixSortStoreAlgorithm>";
   }
+  return "<unknown RadixSortStoreAlgorithm>";
 }
 } // namespace detail
 #endif // _CCCL_HOSTED()
@@ -82,7 +81,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::RadixSortStoreAlgorithm, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::RadixSortStoreAlgorithm, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::RadixSortStoreAlgorithm& algo, FmtCtx& ctx) const
diff --git a/cub/cub/block/block_load.cuh b/cub/cub/block/block_load.cuh
index edd6d340f6a..a6ee301345d 100644
--- a/cub/cub/block/block_load.cuh
+++ b/cub/cub/block/block_load.cuh
@@ -732,7 +732,7 @@ enum BlockLoadAlgorithm
 #if _CCCL_HOSTED() && !defined(_CCCL_DOXYGEN_INVOKED)
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(BlockLoadAlgorithm algo) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(BlockLoadAlgorithm algo) noexcept
 {
   switch (algo)
   {
@@ -748,9 +748,8 @@ namespace detail
       return "BLOCK_LOAD_WARP_TRANSPOSE";
     case BLOCK_LOAD_WARP_TRANSPOSE_TIMESLICED:
       return "BLOCK_LOAD_WARP_TRANSPOSE_TIMESLICED";
-    default:
-      return "<unknown BlockLoadAlgorithm>";
   }
+  return "<unknown BlockLoadAlgorithm>";
 }
 } // namespace detail
 
@@ -764,7 +763,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::BlockLoadAlgorithm, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::BlockLoadAlgorithm, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::BlockLoadAlgorithm& algo, FmtCtx& ctx) const
diff --git a/cub/cub/block/block_radix_rank.cuh b/cub/cub/block/block_radix_rank.cuh
index 7482951c83d..35748d79cd0 100644
--- a/cub/cub/block/block_radix_rank.cuh
+++ b/cub/cub/block/block_radix_rank.cuh
@@ -76,7 +76,7 @@ enum RadixRankAlgorithm
 #if _CCCL_HOSTED() && !defined(_CCCL_DOXYGEN_INVOKED)
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(RadixRankAlgorithm algo) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(RadixRankAlgorithm algo) noexcept
 {
   switch (algo)
   {
@@ -90,9 +90,8 @@ namespace detail
       return "RADIX_RANK_MATCH_EARLY_COUNTS_ANY";
     case RADIX_RANK_MATCH_EARLY_COUNTS_ATOMIC_OR:
       return "RADIX_RANK_MATCH_EARLY_COUNTS_ATOMIC_OR";
-    default:
-      return "<unknown RadixRankAlgorithm>";
   }
+  return "<unknown RadixRankAlgorithm>";
 }
 } // namespace detail
 
@@ -106,7 +105,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::RadixRankAlgorithm, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::RadixRankAlgorithm, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::RadixRankAlgorithm& algo, FmtCtx& ctx) const
diff --git a/cub/cub/block/block_reduce.cuh b/cub/cub/block/block_reduce.cuh
index 6d7da26213b..47ae9976a36 100644
--- a/cub/cub/block/block_reduce.cuh
+++ b/cub/cub/block/block_reduce.cuh
@@ -154,7 +154,7 @@ enum BlockReduceAlgorithm
 #if _CCCL_HOSTED() && !defined(_CCCL_DOXYGEN_INVOKED)
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(BlockReduceAlgorithm algo) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(BlockReduceAlgorithm algo) noexcept
 {
   switch (algo)
   {
@@ -166,9 +166,8 @@ namespace detail
       return "BLOCK_REDUCE_WARP_REDUCTIONS";
     case BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC:
       return "BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC";
-    default:
-      return "<unknown BlockReduceAlgorithm>";
   }
+  return "<unknown BlockReduceAlgorithm>";
 }
 } // namespace detail
 
@@ -182,7 +181,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::BlockReduceAlgorithm, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::BlockReduceAlgorithm, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::BlockReduceAlgorithm& algo, FmtCtx& ctx) const
diff --git a/cub/cub/block/block_scan.cuh b/cub/cub/block/block_scan.cuh
index 195ea5658b9..1f60e1c393d 100644
--- a/cub/cub/block/block_scan.cuh
+++ b/cub/cub/block/block_scan.cuh
@@ -105,7 +105,7 @@ enum BlockScanAlgorithm
 #if _CCCL_HOSTED() && !defined(_CCCL_DOXYGEN_INVOKED)
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(BlockScanAlgorithm algo) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(BlockScanAlgorithm algo) noexcept
 {
   switch (algo)
   {
@@ -115,9 +115,8 @@ namespace detail
       return "BLOCK_SCAN_RAKING_MEMOIZE";
     case BLOCK_SCAN_WARP_SCANS:
       return "BLOCK_SCAN_WARP_SCANS";
-    default:
-      return "<unknown BlockScanAlgorithm>";
   }
+  return "<unknown BlockScanAlgorithm>";
 }
 } // namespace detail
 
@@ -131,7 +130,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::BlockScanAlgorithm, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::BlockScanAlgorithm, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::BlockScanAlgorithm& algo, FmtCtx& ctx) const
diff --git a/cub/cub/block/block_store.cuh b/cub/cub/block/block_store.cuh
index c94dd0bac67..fb82c391e43 100644
--- a/cub/cub/block/block_store.cuh
+++ b/cub/cub/block/block_store.cuh
@@ -548,7 +548,7 @@ enum BlockStoreAlgorithm
 #if _CCCL_HOSTED() && !defined(_CCCL_DOXYGEN_INVOKED)
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(BlockStoreAlgorithm algo) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(BlockStoreAlgorithm algo) noexcept
 {
   switch (algo)
   {
@@ -564,9 +564,8 @@ namespace detail
       return "BLOCK_STORE_WARP_TRANSPOSE";
     case BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED:
       return "BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED";
-    default:
-      return "<unknown BlockStoreAlgorithm>";
   }
+  return "<unknown BlockStoreAlgorithm>";
 }
 } // namespace detail
 
@@ -580,7 +579,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::BlockStoreAlgorithm, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::BlockStoreAlgorithm, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::BlockStoreAlgorithm& algo, FmtCtx& ctx) const
diff --git a/cub/cub/detail/delay_constructor.cuh b/cub/cub/detail/delay_constructor.cuh
index 96fc5d39073..721c09d817d 100644
--- a/cub/cub/detail/delay_constructor.cuh
+++ b/cub/cub/detail/delay_constructor.cuh
@@ -14,110 +14,12 @@
 #endif // no system header
 
 #include <cub/agent/single_pass_scan_operators.cuh>
+#include <cub/device/dispatch/tuning/common.cuh>
 
 #include <cuda/std/__concepts/same_as.h>
-#include <cuda/std/__fwd/format.h>
-#include <cuda/std/__host_stdlib/ostream>
 
 CUB_NAMESPACE_BEGIN
 
-//! The delay algorithm used by decoupled lookback
-enum class LookbackDelayAlgorithm
-{
-  no_delay,
-  fixed_delay,
-  exponential_backoff,
-  exponential_backoff_jitter,
-  exponential_backoff_jitter_window,
-  exponential_backon_jitter_window,
-  exponential_backon_jitter,
-  exponential_backon,
-  __reduce_by_key //!< Internal
-};
-
-#if _CCCL_HOSTED()
-namespace detail
-{
-[[nodiscard]] constexpr const char* to_string(LookbackDelayAlgorithm algo) noexcept
-{
-  switch (algo)
-  {
-    case LookbackDelayAlgorithm::no_delay:
-      return "LookbackDelayAlgorithm::no_delay";
-    case LookbackDelayAlgorithm::fixed_delay:
-      return "LookbackDelayAlgorithm::fixed_delay";
-    case LookbackDelayAlgorithm::exponential_backoff:
-      return "LookbackDelayAlgorithm::exponential_backoff";
-    case LookbackDelayAlgorithm::exponential_backoff_jitter:
-      return "LookbackDelayAlgorithm::exponential_backoff_jitter";
-    case LookbackDelayAlgorithm::exponential_backoff_jitter_window:
-      return "LookbackDelayAlgorithm::exponential_backoff_jitter_window";
-    case LookbackDelayAlgorithm::exponential_backon_jitter_window:
-      return "LookbackDelayAlgorithm::exponential_backon_jitter_window";
-    case LookbackDelayAlgorithm::exponential_backon_jitter:
-      return "LookbackDelayAlgorithm::exponential_backon_jitter";
-    case LookbackDelayAlgorithm::exponential_backon:
-      return "LookbackDelayAlgorithm::exponential_backon";
-    case LookbackDelayAlgorithm::__reduce_by_key:
-      return "LookbackDelayAlgorithm::__reduce_by_key";
-    default:
-      return "<unknown LookbackDelayAlgorithm>";
-  }
-}
-} // namespace detail
-#endif // _CCCL_HOSTED()
-
-#if _CCCL_HOSTED()
-inline ::std::ostream& operator<<(::std::ostream& os, LookbackDelayAlgorithm algo)
-{
-  return os << CUB_NS_QUALIFIER::detail::to_string(algo);
-}
-#endif // _CCCL_HOSTED()
-
-CUB_NAMESPACE_END
-
-#if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
-template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::LookbackDelayAlgorithm, CharT> : formatter<const CharT*, CharT>
-{
-  template <class FmtCtx>
-  auto format(const CUB_NS_QUALIFIER::LookbackDelayAlgorithm& algo, FmtCtx& ctx) const
-  {
-    return formatter<const CharT*, CharT>::format(CUB_NS_QUALIFIER::detail::to_string(algo), ctx);
-  }
-};
-#endif // __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
-
-CUB_NAMESPACE_BEGIN
-
-//! The policy configuring the delay algorithm used by decoupled lookback
-struct LookbackDelayPolicy
-{
-  LookbackDelayAlgorithm kind; //!< The algorithm used for delaying during decoupled lookback
-  unsigned int delay; //!< The delay in nanoseconds
-  unsigned int l2_write_latency; //!< The write latency of the L2 cache in nanoseconds
-
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
-  operator==(const LookbackDelayPolicy& lhs, const LookbackDelayPolicy& rhs) noexcept
-  {
-    return lhs.kind == rhs.kind && lhs.delay == rhs.delay && lhs.l2_write_latency == rhs.l2_write_latency;
-  }
-
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
-  operator!=(const LookbackDelayPolicy& lhs, const LookbackDelayPolicy& rhs) noexcept
-  {
-    return !(lhs == rhs);
-  }
-
-#if _CCCL_HOSTED()
-  friend ::std::ostream& operator<<(::std::ostream& os, const LookbackDelayPolicy& p)
-  {
-    return os << "LookbackDelayPolicy { .kind = " << p.kind << ", .delay = " << p.delay
-              << ", .l2_write_latency = " << p.l2_write_latency << " }";
-  }
-#endif // _CCCL_HOSTED()
-};
-
 namespace detail
 {
 template <typename DelayConstructor>
diff --git a/cub/cub/device/dispatch/tuning/common.cuh b/cub/cub/device/dispatch/tuning/common.cuh
index 9965f3fa2e5..02e524ef464 100644
--- a/cub/cub/device/dispatch/tuning/common.cuh
+++ b/cub/cub/device/dispatch/tuning/common.cuh
@@ -20,7 +20,10 @@
 
 #include <cuda/__functional/maximum.h>
 #include <cuda/__functional/minimum.h>
+#include <cuda/std/__concepts/same_as.h>
 #include <cuda/std/__functional/operations.h>
+#include <cuda/std/__fwd/format.h>
+#include <cuda/std/__host_stdlib/ostream>
 #include <cuda/std/__type_traits/is_signed.h>
 
 CUB_NAMESPACE_BEGIN
@@ -186,4 +189,99 @@ _CCCL_HOST_DEVICE_API constexpr length_size classify_length_size()
   return sizeof(LengthT) == 4 ? length_size::_4 : length_size::unknown;
 }
 } // namespace detail
+
+//! The delay algorithm used by decoupled lookback
+enum class LookbackDelayAlgorithm
+{
+  no_delay,
+  fixed_delay,
+  exponential_backoff,
+  exponential_backoff_jitter,
+  exponential_backoff_jitter_window,
+  exponential_backon_jitter_window,
+  exponential_backon_jitter,
+  exponential_backon,
+  __reduce_by_key //!< Internal
+};
+
+namespace detail
+{
+_CCCL_API [[nodiscard]] constexpr const char* to_string(LookbackDelayAlgorithm algo) noexcept
+{
+  switch (algo)
+  {
+    case LookbackDelayAlgorithm::no_delay:
+      return "LookbackDelayAlgorithm::no_delay";
+    case LookbackDelayAlgorithm::fixed_delay:
+      return "LookbackDelayAlgorithm::fixed_delay";
+    case LookbackDelayAlgorithm::exponential_backoff:
+      return "LookbackDelayAlgorithm::exponential_backoff";
+    case LookbackDelayAlgorithm::exponential_backoff_jitter:
+      return "LookbackDelayAlgorithm::exponential_backoff_jitter";
+    case LookbackDelayAlgorithm::exponential_backoff_jitter_window:
+      return "LookbackDelayAlgorithm::exponential_backoff_jitter_window";
+    case LookbackDelayAlgorithm::exponential_backon_jitter_window:
+      return "LookbackDelayAlgorithm::exponential_backon_jitter_window";
+    case LookbackDelayAlgorithm::exponential_backon_jitter:
+      return "LookbackDelayAlgorithm::exponential_backon_jitter";
+    case LookbackDelayAlgorithm::exponential_backon:
+      return "LookbackDelayAlgorithm::exponential_backon";
+    case LookbackDelayAlgorithm::__reduce_by_key:
+      return "LookbackDelayAlgorithm::__reduce_by_key";
+  }
+  return "<unknown LookbackDelayAlgorithm>";
+}
+} // namespace detail
+
+#if _CCCL_HOSTED()
+inline ::std::ostream& operator<<(::std::ostream& os, LookbackDelayAlgorithm algo)
+{
+  return os << CUB_NS_QUALIFIER::detail::to_string(algo);
+}
+#endif // _CCCL_HOSTED()
+
+CUB_NAMESPACE_END
+
+#if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
+template <::cuda::std::same_as<char> CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::LookbackDelayAlgorithm, CharT> : formatter<const CharT*, CharT>
+{
+  template <class FmtCtx>
+  auto format(const CUB_NS_QUALIFIER::LookbackDelayAlgorithm& algo, FmtCtx& ctx) const
+  {
+    return formatter<const CharT*, CharT>::format(CUB_NS_QUALIFIER::detail::to_string(algo), ctx);
+  }
+};
+#endif // __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
+
+CUB_NAMESPACE_BEGIN
+
+//! The policy configuring the delay algorithm used by decoupled lookback
+struct LookbackDelayPolicy
+{
+  LookbackDelayAlgorithm kind; //!< The algorithm used for delaying during decoupled lookback
+  unsigned int delay; //!< The delay in nanoseconds
+  unsigned int l2_write_latency; //!< The write latency of the L2 cache in nanoseconds
+
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
+  operator==(const LookbackDelayPolicy& lhs, const LookbackDelayPolicy& rhs) noexcept
+  {
+    return lhs.kind == rhs.kind && lhs.delay == rhs.delay && lhs.l2_write_latency == rhs.l2_write_latency;
+  }
+
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
+  operator!=(const LookbackDelayPolicy& lhs, const LookbackDelayPolicy& rhs) noexcept
+  {
+    return !(lhs == rhs);
+  }
+
+#if _CCCL_HOSTED()
+  friend ::std::ostream& operator<<(::std::ostream& os, const LookbackDelayPolicy& p)
+  {
+    return os << "LookbackDelayPolicy { .kind = " << p.kind << ", .delay = " << p.delay
+              << ", .l2_write_latency = " << p.l2_write_latency << " }";
+  }
+#endif // _CCCL_HOSTED()
+};
+
 CUB_NAMESPACE_END
diff --git a/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh b/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh
index a94a19b17b0..cb91fea9be0 100644
--- a/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh
@@ -31,7 +31,7 @@ struct AdjacentDifferencePolicy
   CacheLoadModifier load_modifier; //!< The @ref CacheLoadModifier used for loading items from global memory
   BlockStoreAlgorithm store_algorithm; //!< The @ref BlockStoreAlgorithm used for storing items to global memory
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const AdjacentDifferencePolicy& lhs, const AdjacentDifferencePolicy& rhs)
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -39,7 +39,7 @@ struct AdjacentDifferencePolicy
         && lhs.store_algorithm == rhs.store_algorithm;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const AdjacentDifferencePolicy& lhs, const AdjacentDifferencePolicy& rhs)
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_batch_memcpy.cuh b/cub/cub/device/dispatch/tuning/tuning_batch_memcpy.cuh
index 5dca7d730e7..3bd78f05812 100644
--- a/cub/cub/device/dispatch/tuning/tuning_batch_memcpy.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_batch_memcpy.cuh
@@ -38,7 +38,7 @@ struct BatchedCopySmallBufferPolicy
   LookbackDelayPolicy buffer_lookback_delay; //!< The @ref LookbackDelayPolicy for the buffer offset scan
   LookbackDelayPolicy block_lookback_delay; //!< The @ref LookbackDelayPolicy for the block offset scan
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const BatchedCopySmallBufferPolicy& lhs, const BatchedCopySmallBufferPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.buffers_per_thread == rhs.buffers_per_thread
@@ -50,7 +50,7 @@ struct BatchedCopySmallBufferPolicy
         && lhs.block_lookback_delay == rhs.block_lookback_delay;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const BatchedCopySmallBufferPolicy& lhs, const BatchedCopySmallBufferPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -77,13 +77,13 @@ struct BatchedCopyLargeBufferPolicy
   int threads_per_block; //!< Number of threads in a CUDA block
   int bytes_per_thread; //!< Number of bytes processed per thread
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const BatchedCopyLargeBufferPolicy& lhs, const BatchedCopyLargeBufferPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.bytes_per_thread == rhs.bytes_per_thread;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const BatchedCopyLargeBufferPolicy& lhs, const BatchedCopyLargeBufferPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -104,13 +104,13 @@ struct BatchedCopyPolicy
   BatchedCopySmallBufferPolicy small_buffer; //!< Sub-policy for small buffers copied by a single thread block
   BatchedCopyLargeBufferPolicy large_buffer; //!< Sub-policy for large buffers requiring multi-block collaboration
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const BatchedCopyPolicy& lhs, const BatchedCopyPolicy& rhs) noexcept
   {
     return lhs.small_buffer == rhs.small_buffer && lhs.large_buffer == rhs.large_buffer;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const BatchedCopyPolicy& lhs, const BatchedCopyPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_batched_topk.cuh b/cub/cub/device/dispatch/tuning/tuning_batched_topk.cuh
index 6ea6a2fa7cd..48f976514bd 100644
--- a/cub/cub/device/dispatch/tuning/tuning_batched_topk.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_batched_topk.cuh
@@ -31,13 +31,13 @@ struct epilogue_policy
   BlockStoreAlgorithm store_algorithm;
   BlockScanAlgorithm scan_algorithm;
 
-  _CCCL_HOST_DEVICE_API constexpr friend bool operator==(const epilogue_policy& lhs, const epilogue_policy& rhs)
+  _CCCL_HOST_DEVICE_API friend constexpr bool operator==(const epilogue_policy& lhs, const epilogue_policy& rhs)
   {
     return lhs.items_per_thread == rhs.items_per_thread && lhs.load_algorithm == rhs.load_algorithm
         && lhs.store_algorithm == rhs.store_algorithm && lhs.scan_algorithm == rhs.scan_algorithm;
   }
 
-  _CCCL_HOST_DEVICE_API constexpr friend bool operator!=(const epilogue_policy& lhs, const epilogue_policy& rhs)
+  _CCCL_HOST_DEVICE_API friend constexpr bool operator!=(const epilogue_policy& lhs, const epilogue_policy& rhs)
   {
     return !(lhs == rhs);
   }
@@ -61,14 +61,14 @@ struct worker_policy
 
   epilogue_policy epilogue;
 
-  _CCCL_HOST_DEVICE_API constexpr friend bool operator==(const worker_policy& lhs, const worker_policy& rhs)
+  _CCCL_HOST_DEVICE_API friend constexpr bool operator==(const worker_policy& lhs, const worker_policy& rhs)
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
         && lhs.load_algorithm == rhs.load_algorithm && lhs.store_algorithm == rhs.store_algorithm
         && lhs.epilogue == rhs.epilogue;
   }
 
-  _CCCL_HOST_DEVICE_API constexpr friend bool operator!=(const worker_policy& lhs, const worker_policy& rhs)
+  _CCCL_HOST_DEVICE_API friend constexpr bool operator!=(const worker_policy& lhs, const worker_policy& rhs)
   {
     return !(lhs == rhs);
   }
@@ -88,12 +88,12 @@ struct multi_worker_policy
   int threads_per_block;
   int items_per_thread;
 
-  _CCCL_HOST_DEVICE_API constexpr friend bool operator==(const multi_worker_policy& lhs, const multi_worker_policy& rhs)
+  _CCCL_HOST_DEVICE_API friend constexpr bool operator==(const multi_worker_policy& lhs, const multi_worker_policy& rhs)
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread;
   }
 
-  _CCCL_HOST_DEVICE_API constexpr friend bool operator!=(const multi_worker_policy& lhs, const multi_worker_policy& rhs)
+  _CCCL_HOST_DEVICE_API friend constexpr bool operator!=(const multi_worker_policy& lhs, const multi_worker_policy& rhs)
   {
     return !(lhs == rhs);
   }
@@ -114,13 +114,13 @@ struct batched_topk_policy
   ::cuda::std::array<worker_policy, 6> worker_per_segment_policies;
   multi_worker_policy multi_worker_per_segment_policy;
 
-  _CCCL_HOST_DEVICE_API constexpr friend bool operator==(const batched_topk_policy& lhs, const batched_topk_policy& rhs)
+  _CCCL_HOST_DEVICE_API friend constexpr bool operator==(const batched_topk_policy& lhs, const batched_topk_policy& rhs)
   {
     return lhs.worker_per_segment_policies == rhs.worker_per_segment_policies
         && lhs.multi_worker_per_segment_policy == rhs.multi_worker_per_segment_policy;
   }
 
-  _CCCL_HOST_DEVICE_API constexpr friend bool operator!=(const batched_topk_policy& lhs, const batched_topk_policy& rhs)
+  _CCCL_HOST_DEVICE_API friend constexpr bool operator!=(const batched_topk_policy& lhs, const batched_topk_policy& rhs)
   {
     return !(lhs == rhs);
   }
diff --git a/cub/cub/device/dispatch/tuning/tuning_find.cuh b/cub/cub/device/dispatch/tuning/tuning_find.cuh
index f89618831e5..14fff9063ec 100644
--- a/cub/cub/device/dispatch/tuning/tuning_find.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_find.cuh
@@ -30,14 +30,14 @@ struct FindIfPolicy
   int vec_size; //!< Vectorization size for loading items
   CacheLoadModifier load_modifier; //!< The @ref CacheLoadModifier used for loading items from global memory
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const FindIfPolicy& lhs, const FindIfPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
         && lhs.vec_size == rhs.vec_size && lhs.load_modifier == rhs.load_modifier;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const FindIfPolicy& lhs, const FindIfPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_find_bound_sorted_values.cuh b/cub/cub/device/dispatch/tuning/tuning_find_bound_sorted_values.cuh
index cc849e45a86..badfc6c1abf 100644
--- a/cub/cub/device/dispatch/tuning/tuning_find_bound_sorted_values.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_find_bound_sorted_values.cuh
@@ -33,14 +33,14 @@ struct FindBoundSortedValuesPolicy
   int items_per_thread; //!< Number of items processed per thread
   CacheLoadModifier load_modifier; //!< The @ref CacheLoadModifier used for loading items from global memory
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const FindBoundSortedValuesPolicy& lhs, const FindBoundSortedValuesPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
         && lhs.load_modifier == rhs.load_modifier;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const FindBoundSortedValuesPolicy& lhs, const FindBoundSortedValuesPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_for.cuh b/cub/cub/device/dispatch/tuning/tuning_for.cuh
index e0537dd4f40..9f1de9703ca 100644
--- a/cub/cub/device/dispatch/tuning/tuning_for.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_for.cuh
@@ -27,13 +27,13 @@ struct ForPolicy
                          //!< runtime based on the maximum occupancy of the kernel.
   int items_per_thread; //!< Number of items processed per thread
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const ForPolicy& lhs, const ForPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const ForPolicy& lhs, const ForPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh
index 20bdf1578b1..c3e2e73c0ce 100644
--- a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh
@@ -40,7 +40,7 @@ struct HistogramPolicy
   int init_kernel_pdl_trigger_max_bins; //!< Maximum number of bins for the init kernel to trigger the histogram kernel
                                         //!< early using PDL
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const HistogramPolicy& lhs, const HistogramPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.pixels_per_thread == rhs.pixels_per_thread
@@ -50,7 +50,7 @@ struct HistogramPolicy
         && lhs.init_kernel_pdl_trigger_max_bins == rhs.init_kernel_pdl_trigger_max_bins;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const HistogramPolicy& lhs, const HistogramPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_merge.cuh b/cub/cub/device/dispatch/tuning/tuning_merge.cuh
index 433d68ae62d..e44a9bdbe40 100644
--- a/cub/cub/device/dispatch/tuning/tuning_merge.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_merge.cuh
@@ -40,7 +40,7 @@ struct MergePolicy
   bool use_bulk_copy_for_values; //!< Whether to use bulk copy (cp.async.bulk) for loading values into shared memory
   bool unroll = true; //<! Whether to unroll the loops inside the serial merge implementation
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const MergePolicy& lhs, const MergePolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -49,7 +49,7 @@ struct MergePolicy
         && lhs.use_bulk_copy_for_values == rhs.use_bulk_copy_for_values && lhs.unroll == rhs.unroll;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const MergePolicy& lhs, const MergePolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh
index 94928b0de4f..db8def1bb83 100644
--- a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh
@@ -69,7 +69,7 @@ struct MergeSortPolicy
   bool unroll = true;
 #endif // CCCL_AVOID_SORT_UNROLL
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const MergeSortPolicy& lhs, const MergeSortPolicy& rhs)
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -77,7 +77,7 @@ struct MergeSortPolicy
         && lhs.store_algorithm == rhs.store_algorithm && lhs.unroll == rhs.unroll;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const MergeSortPolicy& lhs, const MergeSortPolicy& rhs)
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh
index dae9bd75e6e..3b248156f3c 100644
--- a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh
@@ -41,7 +41,7 @@ enum class RadixSortAlgorithm
 #if _CCCL_HOSTED()
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(RadixSortAlgorithm algo) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(RadixSortAlgorithm algo) noexcept
 {
   switch (algo)
   {
@@ -49,9 +49,8 @@ namespace detail
       return "RadixSortAlgorithm::multi_pass";
     case RadixSortAlgorithm::onesweep:
       return "RadixSortAlgorithm::onesweep";
-    default:
-      return "<unknown RadixSortAlgorithm>";
   }
+  return "<unknown RadixSortAlgorithm>";
 }
 } // namespace detail
 #endif // _CCCL_HOSTED()
@@ -67,7 +66,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::RadixSortAlgorithm, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::RadixSortAlgorithm, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::RadixSortAlgorithm& algo, FmtCtx& ctx) const
@@ -90,14 +89,14 @@ struct RadixSortHistogramPolicy
   int private_partitions;
   int radix_bits; //!< Number of bits per radix digit
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const RadixSortHistogramPolicy& lhs, const RadixSortHistogramPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
         && lhs.private_partitions == rhs.private_partitions && lhs.radix_bits == rhs.radix_bits;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const RadixSortHistogramPolicy& lhs, const RadixSortHistogramPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -119,13 +118,13 @@ struct RadixSortExclusiveSumPolicy
   int threads_per_block; //!< Number of threads in a CUDA block
   int radix_bits; //!< Number of bits per radix digit
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const RadixSortExclusiveSumPolicy& lhs, const RadixSortExclusiveSumPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.radix_bits == rhs.radix_bits;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const RadixSortExclusiveSumPolicy& lhs, const RadixSortExclusiveSumPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -156,7 +155,7 @@ struct RadixSortOnesweepPolicy
 
   int radix_bits; //!< Number of bits per radix digit
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const RadixSortOnesweepPolicy& lhs, const RadixSortOnesweepPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -165,7 +164,7 @@ struct RadixSortOnesweepPolicy
         && lhs.radix_bits == rhs.radix_bits;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const RadixSortOnesweepPolicy& lhs, const RadixSortOnesweepPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -194,7 +193,7 @@ struct RadixSortDownsweepPolicy
   BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for scanning within a thread block
   int radix_bits; //!< Number of bits per radix digit
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const RadixSortDownsweepPolicy& lhs, const RadixSortDownsweepPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -203,7 +202,7 @@ struct RadixSortDownsweepPolicy
         && lhs.radix_bits == rhs.radix_bits;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const RadixSortDownsweepPolicy& lhs, const RadixSortDownsweepPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -229,14 +228,14 @@ struct RadixSortUpsweepPolicy
   CacheLoadModifier load_modifier; //!< The @ref CacheLoadModifier used for loading items from global memory
   int radix_bits; //!< Number of bits per radix digit
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const RadixSortUpsweepPolicy& lhs, const RadixSortUpsweepPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
         && lhs.load_modifier == rhs.load_modifier && lhs.radix_bits == rhs.radix_bits;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const RadixSortUpsweepPolicy& lhs, const RadixSortUpsweepPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -266,7 +265,7 @@ struct RadixSortPolicy
   RadixSortUpsweepPolicy alt_upsweep; //!< Alternate upsweep pass policy with fewer radix bits
   RadixSortDownsweepPolicy single_tile; //!< Single-tile sort policy for small inputs
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const RadixSortPolicy& lhs, const RadixSortPolicy& rhs) noexcept
   {
     return lhs.algorithm == rhs.algorithm && lhs.histogram == rhs.histogram && lhs.exclusive_sum == rhs.exclusive_sum
@@ -275,7 +274,7 @@ struct RadixSortPolicy
         && lhs.single_tile == rhs.single_tile;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const RadixSortPolicy& lhs, const RadixSortPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_reduce.cuh b/cub/cub/device/dispatch/tuning/tuning_reduce.cuh
index c2f262943ae..0819aca15de 100644
--- a/cub/cub/device/dispatch/tuning/tuning_reduce.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_reduce.cuh
@@ -35,7 +35,7 @@ struct ReducePassPolicy
   BlockReduceAlgorithm reduce_algorithm; //!< The @ref BlockReduceAlgorithm to use
   CacheLoadModifier load_modifier; //!< The @ref CacheLoadModifier used for loading items from global memory
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const ReducePassPolicy& lhs, const ReducePassPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -43,7 +43,7 @@ struct ReducePassPolicy
         && lhs.load_modifier == rhs.load_modifier;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const ReducePassPolicy& lhs, const ReducePassPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -66,13 +66,13 @@ struct ReducePolicy
   ReducePassPolicy single_tile; //!< Policy used for the single-tile pass. Used as second pass after the multi-tile pass
                                 //!< in some cases, or when the problem size fits into a single tile.
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const ReducePolicy& lhs, const ReducePolicy& rhs) noexcept
   {
     return lhs.multi_tile == rhs.multi_tile && lhs.single_tile == rhs.single_tile;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const ReducePolicy& lhs, const ReducePolicy& rhs) noexcept
   {
     return !(lhs == rhs);
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..6672f25dea4 100644
--- a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh
@@ -41,7 +41,7 @@ struct ReduceByKeyPolicy
   BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for the prefix scan
   LookbackDelayPolicy lookback_delay; //!< The @ref LookbackDelayPolicy used for the lookback delay
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const ReduceByKeyPolicy& lhs, const ReduceByKeyPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -49,7 +49,7 @@ struct ReduceByKeyPolicy
         && lhs.scan_algorithm == rhs.scan_algorithm && lhs.lookback_delay == rhs.lookback_delay;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const ReduceByKeyPolicy& lhs, const ReduceByKeyPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh b/cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh
index 30b3e343a95..2491da7228e 100644
--- a/cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh
@@ -43,7 +43,7 @@ struct RleEncodePolicy
   BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for the prefix scan
   LookbackDelayPolicy lookback_delay; //!< The @ref LookbackDelayPolicy used for the lookback delay
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const RleEncodePolicy& lhs, const RleEncodePolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -51,7 +51,7 @@ struct RleEncodePolicy
         && lhs.scan_algorithm == rhs.scan_algorithm && lhs.lookback_delay == rhs.lookback_delay;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const RleEncodePolicy& lhs, const RleEncodePolicy& rhs) noexcept
   {
     return !(lhs == rhs);
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..2159451606c 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
@@ -44,7 +44,7 @@ struct RleNonTrivialRunsPolicy
   //! The @ref LookbackDelayPolicy controlling the delay between lookback iterations
   LookbackDelayPolicy lookback_delay = {LookbackDelayAlgorithm::fixed_delay, 350, 450};
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const RleNonTrivialRunsPolicy& lhs, const RleNonTrivialRunsPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -53,7 +53,7 @@ struct RleNonTrivialRunsPolicy
         && lhs.lookback_delay == rhs.lookback_delay;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const RleNonTrivialRunsPolicy& lhs, const RleNonTrivialRunsPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_scan.cuh b/cub/cub/device/dispatch/tuning/tuning_scan.cuh
index 3e9204a036e..709732ed726 100644
--- a/cub/cub/device/dispatch/tuning/tuning_scan.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_scan.cuh
@@ -55,7 +55,7 @@ enum class ScanAlgorithm
 #if _CCCL_HOSTED()
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(ScanAlgorithm algo) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(ScanAlgorithm algo) noexcept
 {
   switch (algo)
   {
@@ -63,9 +63,8 @@ namespace detail
       return "ScanAlgorithm::lookback";
     case ScanAlgorithm::lookahead:
       return "ScanAlgorithm::lookahead";
-    default:
-      return "<unknown ScanAlgorithm>";
   }
+  return "<unknown ScanAlgorithm>";
 }
 } // namespace detail
 #endif // _CCCL_HOSTED()
@@ -81,7 +80,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::ScanAlgorithm, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::ScanAlgorithm, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::ScanAlgorithm& algo, FmtCtx& ctx) const
@@ -105,7 +104,7 @@ struct ScanLookbackPolicy
   BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for scanning within a thread block
   LookbackDelayPolicy lookback_delay; //!< The policy configuring the delay used in decoupled lookback
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const ScanLookbackPolicy& lhs, const ScanLookbackPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -114,7 +113,7 @@ struct ScanLookbackPolicy
         && lhs.lookback_delay == rhs.lookback_delay;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const ScanLookbackPolicy& lhs, const ScanLookbackPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -156,7 +155,7 @@ struct ScanLookaheadPolicy
     return items_per_thread * reduce_and_scan_warps * cub::detail::warp_threads;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const ScanLookaheadPolicy& lhs, const ScanLookaheadPolicy& rhs) noexcept
   {
     return lhs.reduce_and_scan_warps == rhs.reduce_and_scan_warps && lhs.items_per_thread == rhs.items_per_thread
@@ -164,7 +163,7 @@ struct ScanLookaheadPolicy
         && lhs.lookahead_stages == rhs.lookahead_stages && lhs.block_idx_stages == rhs.block_idx_stages;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const ScanLookaheadPolicy& lhs, const ScanLookaheadPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -188,12 +187,12 @@ struct ScanPolicy
   ScanLookbackPolicy lookback; //!< The look-back scan policy (used when algorithm is @p lookback, otherwise ignored)
   ScanLookaheadPolicy lookahead; //!< The lookahead scan policy (used when algorithm is @p lookahead, otherwise ignored)
 
-  [[nodiscard]] _CCCL_API constexpr friend bool operator==(const ScanPolicy& lhs, const ScanPolicy& rhs) noexcept
+  [[nodiscard]] _CCCL_API friend constexpr bool operator==(const ScanPolicy& lhs, const ScanPolicy& rhs) noexcept
   {
     return lhs.lookback == rhs.lookback && lhs.lookahead == rhs.lookahead && lhs.algorithm == rhs.algorithm;
   }
 
-  [[nodiscard]] _CCCL_API constexpr friend bool operator!=(const ScanPolicy& lhs, const ScanPolicy& rhs) noexcept
+  [[nodiscard]] _CCCL_API friend constexpr bool operator!=(const ScanPolicy& lhs, const ScanPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
   }
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..95922e9e2b3 100644
--- a/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh
@@ -43,7 +43,7 @@ struct ScanByKeyPolicy
   BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for scanning within a thread block
   LookbackDelayPolicy lookback_delay; //!< The policy configuring the delay used in decoupled lookback
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const ScanByKeyPolicy& lhs, const ScanByKeyPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -52,7 +52,7 @@ struct ScanByKeyPolicy
         && lhs.lookback_delay == rhs.lookback_delay;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const ScanByKeyPolicy& lhs, const ScanByKeyPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_segmented_radix_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_segmented_radix_sort.cuh
index a7823616c15..e70ad62f8f1 100644
--- a/cub/cub/device/dispatch/tuning/tuning_segmented_radix_sort.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_segmented_radix_sort.cuh
@@ -26,13 +26,13 @@ struct SegmentedRadixSortPolicy
   RadixSortDownsweepPolicy regular_pass; //!< Policy for the regular radix sort pass on each segment
   RadixSortDownsweepPolicy alternate_pass; //!< Policy for the alternate radix sort pass on each segment
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const SegmentedRadixSortPolicy& lhs, const SegmentedRadixSortPolicy& rhs) noexcept
   {
     return lhs.regular_pass == rhs.regular_pass && lhs.alternate_pass == rhs.alternate_pass;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const SegmentedRadixSortPolicy& lhs, const SegmentedRadixSortPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh b/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh
index a1cf9919c59..86cf977ef0f 100644
--- a/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh
@@ -38,7 +38,7 @@ struct SegmentedReduceWarpReducePolicy
     return threads_per_block / threads_per_warp;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const SegmentedReduceWarpReducePolicy& lhs, const SegmentedReduceWarpReducePolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.threads_per_warp == rhs.threads_per_warp
@@ -46,7 +46,7 @@ struct SegmentedReduceWarpReducePolicy
         && lhs.load_modifier == rhs.load_modifier;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const SegmentedReduceWarpReducePolicy& lhs, const SegmentedReduceWarpReducePolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -69,14 +69,14 @@ struct SegmentedReducePolicy
   SegmentedReduceWarpReducePolicy medium_reduce; //!< Policy used for medium segments (one warp per segment)
   SegmentedReduceWarpReducePolicy small_reduce; //!< Policy used for small segments (one thread per segment)
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const SegmentedReducePolicy& lhs, const SegmentedReducePolicy& rhs) noexcept
   {
     return lhs.large_reduce == rhs.large_reduce && lhs.medium_reduce == rhs.medium_reduce
         && lhs.small_reduce == rhs.small_reduce;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const SegmentedReducePolicy& lhs, const SegmentedReducePolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_segmented_scan.cuh b/cub/cub/device/dispatch/tuning/tuning_segmented_scan.cuh
index cd4a7bf7f19..e0cbc7f78ac 100644
--- a/cub/cub/device/dispatch/tuning/tuning_segmented_scan.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_segmented_scan.cuh
@@ -36,7 +36,7 @@ struct SegmentedScanBlockPolicy
   BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for block scanning
   int max_segments; //!< Maximum number of segments processed per block
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const SegmentedScanBlockPolicy& lhs, const SegmentedScanBlockPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -45,7 +45,7 @@ struct SegmentedScanBlockPolicy
         && lhs.max_segments == rhs.max_segments;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const SegmentedScanBlockPolicy& lhs, const SegmentedScanBlockPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -69,13 +69,13 @@ struct SegmentedScanPolicy
 {
   SegmentedScanBlockPolicy block; //!< Policy for the block-level segmented scan kernel
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const SegmentedScanPolicy& lhs, const SegmentedScanPolicy& rhs) noexcept
   {
     return lhs.block == rhs.block;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const SegmentedScanPolicy& lhs, const SegmentedScanPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh
index 36a504984fa..2cddb2e3356 100644
--- a/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh
@@ -37,7 +37,7 @@ struct SegmentedSortRadixSortPolicy
   BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for the internal digit-count scan
   int radix_bits; //!< Number of bits per radix digit pass
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const SegmentedSortRadixSortPolicy& lhs, const SegmentedSortRadixSortPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -46,7 +46,7 @@ struct SegmentedSortRadixSortPolicy
         && lhs.radix_bits == rhs.radix_bits;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const SegmentedSortRadixSortPolicy& lhs, const SegmentedSortRadixSortPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -84,7 +84,7 @@ struct SegmentedSortSubWarpMergeSortPolicy
     return threads_per_warp * items_per_thread;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const SegmentedSortSubWarpMergeSortPolicy& lhs, const SegmentedSortSubWarpMergeSortPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.threads_per_warp == rhs.threads_per_warp
@@ -92,7 +92,7 @@ struct SegmentedSortSubWarpMergeSortPolicy
         && lhs.load_modifier == rhs.load_modifier && lhs.store_algorithm == rhs.store_algorithm;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const SegmentedSortSubWarpMergeSortPolicy& lhs, const SegmentedSortSubWarpMergeSortPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -119,14 +119,14 @@ struct SegmentedSortPolicy
   int partitioning_threshold; //!< Number of segments above which different algorithms will be used for different size
                               //!< buckets
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const SegmentedSortPolicy& lhs, const SegmentedSortPolicy& rhs) noexcept
   {
     return lhs.large_segment == rhs.large_segment && lhs.small_segment == rhs.small_segment
         && lhs.medium_segment == rhs.medium_segment && lhs.partitioning_threshold == rhs.partitioning_threshold;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const SegmentedSortPolicy& lhs, const SegmentedSortPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_select_if.cuh b/cub/cub/device/dispatch/tuning/tuning_select_if.cuh
index b80d4e6d684..2dadcbdbfc4 100644
--- a/cub/cub/device/dispatch/tuning/tuning_select_if.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_select_if.cuh
@@ -44,7 +44,7 @@ struct SelectPolicy
   BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for scanning
   LookbackDelayPolicy lookback_delay; //!< The policy configuring the delay used in decoupled lookback
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const SelectPolicy& lhs, const SelectPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -52,7 +52,7 @@ struct SelectPolicy
         && lhs.scan_algorithm == rhs.scan_algorithm && lhs.lookback_delay == rhs.lookback_delay;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const SelectPolicy& lhs, const SelectPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -81,7 +81,7 @@ struct PartitionPolicy
   BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for scanning
   LookbackDelayPolicy lookback_delay; //!< The policy configuring the delay used in decoupled lookback
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const PartitionPolicy& lhs, const PartitionPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -89,7 +89,7 @@ struct PartitionPolicy
         && lhs.scan_algorithm == rhs.scan_algorithm && lhs.lookback_delay == rhs.lookback_delay;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const PartitionPolicy& lhs, const PartitionPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
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..16a5d28a0fc 100644
--- a/cub/cub/device/dispatch/tuning/tuning_three_way_partition.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_three_way_partition.cuh
@@ -39,7 +39,7 @@ struct ThreeWayPartitionPolicy
   BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for scanning
   LookbackDelayPolicy lookback_delay; //!< The @ref LookbackDelayPolicy configuring the delay used in decoupled lookback
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const ThreeWayPartitionPolicy& lhs, const ThreeWayPartitionPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -47,7 +47,7 @@ struct ThreeWayPartitionPolicy
         && lhs.scan_algorithm == rhs.scan_algorithm && lhs.lookback_delay == rhs.lookback_delay;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const ThreeWayPartitionPolicy& lhs, const ThreeWayPartitionPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/device/dispatch/tuning/tuning_topk.cuh b/cub/cub/device/dispatch/tuning/tuning_topk.cuh
index 7c36b9c1b97..4d70e70acce 100644
--- a/cub/cub/device/dispatch/tuning/tuning_topk.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_topk.cuh
@@ -54,14 +54,14 @@ struct topk_policy
   BlockScanAlgorithm scan_algorithm;
   int bits_per_pass;
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool operator==(const topk_policy& lhs, const topk_policy& rhs)
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool operator==(const topk_policy& lhs, const topk_policy& rhs)
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
         && lhs.load_algorithm == rhs.load_algorithm && lhs.scan_algorithm == rhs.scan_algorithm
         && lhs.bits_per_pass == rhs.bits_per_pass;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool operator!=(const topk_policy& lhs, const topk_policy& rhs)
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool operator!=(const topk_policy& lhs, const topk_policy& rhs)
   {
     return !(lhs == rhs);
   }
diff --git a/cub/cub/device/dispatch/tuning/tuning_transform.cuh b/cub/cub/device/dispatch/tuning/tuning_transform.cuh
index 2d5eb6bb57b..7ce6d997a53 100644
--- a/cub/cub/device/dispatch/tuning/tuning_transform.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_transform.cuh
@@ -33,7 +33,7 @@
 
 CUB_NAMESPACE_BEGIN
 
-//! Backend algorithms for @ref DeviceTransform.
+//! Backend algorithms for @ref cub::DeviceTransform "DeviceTransform".
 enum class TransformAlgorithm
 {
   // We previously had a fallback algorithm that would use cub::DeviceFor. Benchmarks showed that the prefetch algorithm
@@ -47,7 +47,7 @@ enum class TransformAlgorithm
 #if _CCCL_HOSTED()
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(TransformAlgorithm algo) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(TransformAlgorithm algo) noexcept
 {
   switch (algo)
   {
@@ -59,9 +59,8 @@ namespace detail
       return "TransformAlgorithm::ldgsts";
     case TransformAlgorithm::ublkcp:
       return "TransformAlgorithm::ublkcp";
-    default:
-      return "<unknown TransformAlgorithm>";
   }
+  return "<unknown TransformAlgorithm>";
 }
 } // namespace detail
 #endif // _CCCL_HOSTED()
@@ -77,7 +76,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::TransformAlgorithm, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::TransformAlgorithm, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::TransformAlgorithm& algo, FmtCtx& ctx) const
@@ -106,7 +105,7 @@ struct TransformPrefetchPolicy
                          //!< retains the compiler's default unrolling by specifying no unroll pragma. 1 prevents
                          //!< unrolling.
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const TransformPrefetchPolicy& lhs, const TransformPrefetchPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block
@@ -115,7 +114,7 @@ struct TransformPrefetchPolicy
         && lhs.prefetch_byte_stride == rhs.prefetch_byte_stride && lhs.unroll_factor == rhs.unroll_factor;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const TransformPrefetchPolicy& lhs, const TransformPrefetchPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -141,14 +140,14 @@ struct TransformVectorizedPolicy
   int items_per_thread; //!< Number of items processed per thread. Must be a multiple of vec_size.
   int vec_size; //!< Number of elements loaded/stored per vectorized access. Must evenly divide items_per_thread.
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const TransformVectorizedPolicy& lhs, const TransformVectorizedPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
         && lhs.vec_size == rhs.vec_size;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const TransformVectorizedPolicy& lhs, const TransformVectorizedPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -181,7 +180,7 @@ struct TransformAsyncCopyPolicy
                           //!< (16 / sizeof(output), a 16-byte STG.128); 1 disables vectorization (scalar stores) and
                           //!< compiles the vectorized branch out of the kernel.
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const TransformAsyncCopyPolicy& lhs, const TransformAsyncCopyPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.min_items_per_thread == rhs.min_items_per_thread
@@ -189,7 +188,7 @@ struct TransformAsyncCopyPolicy
         && lhs.store_vec_size == rhs.store_vec_size;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const TransformAsyncCopyPolicy& lhs, const TransformAsyncCopyPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
@@ -207,7 +206,7 @@ struct TransformAsyncCopyPolicy
 #endif // _CCCL_HOSTED()
 };
 
-//! The tuning policy for all algorithms in @ref DeviceTransform.
+//! The tuning policy for all algorithms in @ref cub::DeviceTransform "DeviceTransform".
 struct TransformPolicy
 {
   int min_bytes_in_flight; //!< Minimum number of bytes in flight per SM to reach by scaling the items per thread. Has
@@ -221,14 +220,14 @@ struct TransformPolicy
   TransformAsyncCopyPolicy async_copy; //!< Sub-policy for the async copy algorithms. Only used when @p algorithm is @p
                                        //!< ldgsts or @p ublkcp.
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const TransformPolicy& lhs, const TransformPolicy& rhs) noexcept
   {
     return lhs.min_bytes_in_flight == rhs.min_bytes_in_flight && lhs.algorithm == rhs.algorithm
         && lhs.prefetch == rhs.prefetch && lhs.vectorized == rhs.vectorized && lhs.async_copy == rhs.async_copy;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const TransformPolicy& lhs, const TransformPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
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..2f95df779cb 100644
--- a/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh
@@ -39,7 +39,7 @@ struct UniqueByKeyPolicy
   BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for scanning
   LookbackDelayPolicy lookback_delay; //!< The policy configuring the delay used in decoupled lookback
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator==(const UniqueByKeyPolicy& lhs, const UniqueByKeyPolicy& rhs) noexcept
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
@@ -47,7 +47,7 @@ struct UniqueByKeyPolicy
         && lhs.scan_algorithm == rhs.scan_algorithm && lhs.lookback_delay == rhs.lookback_delay;
   }
 
-  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool
   operator!=(const UniqueByKeyPolicy& lhs, const UniqueByKeyPolicy& rhs) noexcept
   {
     return !(lhs == rhs);
diff --git a/cub/cub/thread/thread_load.cuh b/cub/cub/thread/thread_load.cuh
index a881322e1b2..3b86cad7b53 100644
--- a/cub/cub/thread/thread_load.cuh
+++ b/cub/cub/thread/thread_load.cuh
@@ -51,7 +51,7 @@ enum CacheLoadModifier
 #if _CCCL_HOSTED()
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(CacheLoadModifier modifier) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(CacheLoadModifier modifier) noexcept
 {
   switch (modifier)
   {
@@ -69,9 +69,8 @@ namespace detail
       return "LOAD_LDG";
     case LOAD_VOLATILE:
       return "LOAD_VOLATILE";
-    default:
-      return "<unknown CacheLoadModifier>";
   }
+  return "<unknown CacheLoadModifier>";
 }
 } // namespace detail
 #endif // _CCCL_HOSTED()
@@ -87,7 +86,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::CacheLoadModifier, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::CacheLoadModifier, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::CacheLoadModifier& modifier, FmtCtx& ctx) const
diff --git a/cub/cub/warp/warp_load.cuh b/cub/cub/warp/warp_load.cuh
index 3ecdc9d5e2f..de3c2a46d9b 100644
--- a/cub/cub/warp/warp_load.cuh
+++ b/cub/cub/warp/warp_load.cuh
@@ -112,7 +112,7 @@ enum WarpLoadAlgorithm
 #if _CCCL_HOSTED()
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(WarpLoadAlgorithm algo) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(WarpLoadAlgorithm algo) noexcept
 {
   switch (algo)
   {
@@ -124,9 +124,8 @@ namespace detail
       return "WARP_LOAD_VECTORIZE";
     case WARP_LOAD_TRANSPOSE:
       return "WARP_LOAD_TRANSPOSE";
-    default:
-      return "<unknown WarpLoadAlgorithm>";
   }
+  return "<unknown WarpLoadAlgorithm>";
 }
 } // namespace detail
 #endif // _CCCL_HOSTED()
@@ -142,7 +141,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::WarpLoadAlgorithm, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::WarpLoadAlgorithm, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::WarpLoadAlgorithm& algo, FmtCtx& ctx) const
diff --git a/cub/cub/warp/warp_store.cuh b/cub/cub/warp/warp_store.cuh
index 308e9c6aaef..adbf41bbd8b 100644
--- a/cub/cub/warp/warp_store.cuh
+++ b/cub/cub/warp/warp_store.cuh
@@ -116,7 +116,7 @@ enum WarpStoreAlgorithm
 #if _CCCL_HOSTED()
 namespace detail
 {
-[[nodiscard]] constexpr const char* to_string(WarpStoreAlgorithm algo) noexcept
+_CCCL_API [[nodiscard]] constexpr const char* to_string(WarpStoreAlgorithm algo) noexcept
 {
   switch (algo)
   {
@@ -128,9 +128,8 @@ namespace detail
       return "WARP_STORE_VECTORIZE";
     case WARP_STORE_TRANSPOSE:
       return "WARP_STORE_TRANSPOSE";
-    default:
-      return "<unknown WarpStoreAlgorithm>";
   }
+  return "<unknown WarpStoreAlgorithm>";
 }
 } // namespace detail
 
@@ -144,7 +143,7 @@ CUB_NAMESPACE_END
 
 #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED)
 template <::cuda::std::same_as<char> CharT>
-struct std::formatter<CUB_NS_QUALIFIER::WarpStoreAlgorithm, CharT> : formatter<const CharT*, CharT>
+struct ::std::formatter<CUB_NS_QUALIFIER::WarpStoreAlgorithm, CharT> : formatter<const CharT*, CharT>
 {
   template <class FmtCtx>
   auto format(const CUB_NS_QUALIFIER::WarpStoreAlgorithm& algo, FmtCtx& ctx) const
diff --git a/docs/cub/Doxyfile b/docs/cub/Doxyfile
index 6f2dec25329..4d677c3cea4 100644
--- a/docs/cub/Doxyfile
+++ b/docs/cub/Doxyfile
@@ -14,11 +14,12 @@ INPUT                  = ../../cub/cub \
                          ../../cub/cub/warp \
                          ../../cub/cub/block \
                          ../../cub/cub/device \
+                         ../../cub/cub/device/dispatch/tuning \
                          ../../cub/cub/grid \
                          ../../cub/cub/iterator
 
 RECURSIVE              = YES
-EXCLUDE_PATTERNS       = */detail/* */dispatch/* */kernels/* */test/* */examples/*
+EXCLUDE_PATTERNS       = */detail/* */dispatch/dispatch_* */dispatch/kernels/* */kernels/* */test/* */examples/*
 EXCLUDE_SYMBOLS        = *detail* CUB_DETAIL*
 
 FILE_PATTERNS          = *.cuh *.h
@@ -169,6 +170,7 @@ PREDEFINED             = __device__= \
                         "THRUST_NAMESPACE_BEGIN=namespace thrust {" \
                         "THRUST_NAMESPACE_END=}" \
                         "THRUST_PREVENT_MACRO_SUBSTITUTION" \
+                        "_CCCL_HOSTED()=1" \
                         _CCCL_DOXYGEN_INVOKED
 
 # Quiet mode
diff --git a/docs/cudax/Doxyfile b/docs/cudax/Doxyfile
index 4c4bd3320ff..f48e46958e4 100644
--- a/docs/cudax/Doxyfile
+++ b/docs/cudax/Doxyfile
@@ -154,7 +154,8 @@ PREDEFINED = \
     "THRUST_FWD(x)=x" \
     "THRUST_NAMESPACE_BEGIN=namespace thrust {" \
     "THRUST_NAMESPACE_END=}" \
-    "THRUST_PREVENT_MACRO_SUBSTITUTION"
+    "THRUST_PREVENT_MACRO_SUBSTITUTION" \
+    "_CCCL_HOSTED()=1"
 
 # Additional settings from repo.toml
 DISTRIBUTE_GROUP_DOC   = YES
diff --git a/docs/libcudacxx/Doxyfile b/docs/libcudacxx/Doxyfile
index b4a4622186c..8804bb269dc 100644
--- a/docs/libcudacxx/Doxyfile
+++ b/docs/libcudacxx/Doxyfile
@@ -161,7 +161,8 @@ PREDEFINED = \
     "THRUST_FWD(x)=x" \
     "THRUST_NAMESPACE_BEGIN=namespace thrust {" \
     "THRUST_NAMESPACE_END=}" \
-    "THRUST_PREVENT_MACRO_SUBSTITUTION"
+    "THRUST_PREVENT_MACRO_SUBSTITUTION" \
+    "_CCCL_HOSTED()=1"
 
 # IMPORTANT: Aliases for custom commands
 # The rst alias enables embedding reStructuredText in doxygen comments
diff --git a/docs/thrust/Doxyfile b/docs/thrust/Doxyfile
index 7ba9121f78f..0468c283d2d 100644
--- a/docs/thrust/Doxyfile
+++ b/docs/thrust/Doxyfile
@@ -149,7 +149,8 @@ PREDEFINED             = __device__= \
                         "CUB_RDC_ENABLED" \
                         "CUB_NAMESPACE_BEGIN=namespace cub {" \
                         "CUB_NAMESPACE_END=}" \
-                        "CUB_RUNTIME_FUNCTION="
+                        "CUB_RUNTIME_FUNCTION=" \
+                        "_CCCL_HOSTED()=1"
 
 # Images and examples
 IMAGE_PATH             = ../img