From 61942fc6100336c7095c6ede3f3b58f45e67fa0e Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 8 Jul 2026 10:38:36 +0200 Subject: [PATCH 1/8] Move delay constructor policies to common tuning header --- cub/cub/detail/delay_constructor.cuh | 100 +--------------------- cub/cub/device/dispatch/tuning/common.cuh | 100 ++++++++++++++++++++++ 2 files changed, 101 insertions(+), 99 deletions(-) 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 +#include #include -#include -#include 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 ""; - } -} -} // 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 CharT> -struct std::formatter : formatter -{ - template - auto format(const CUB_NS_QUALIFIER::LookbackDelayAlgorithm& algo, FmtCtx& ctx) const - { - return formatter::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 diff --git a/cub/cub/device/dispatch/tuning/common.cuh b/cub/cub/device/dispatch/tuning/common.cuh index 9965f3fa2e5..c7640a1abb9 100644 --- a/cub/cub/device/dispatch/tuning/common.cuh +++ b/cub/cub/device/dispatch/tuning/common.cuh @@ -21,6 +21,8 @@ #include #include #include +#include +#include #include CUB_NAMESPACE_BEGIN @@ -186,4 +188,102 @@ _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 +}; + +#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 ""; + } +} +} // 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 CharT> +struct std::formatter : formatter +{ + template + auto format(const CUB_NS_QUALIFIER::LookbackDelayAlgorithm& algo, FmtCtx& ctx) const + { + return formatter::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() +}; + CUB_NAMESPACE_END From e419fef99730db61aaae94ae7b38dd9078d136cd Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 8 Jul 2026 10:38:49 +0200 Subject: [PATCH 2/8] Include tuning headers for Doxygen --- docs/cub/Doxyfile | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/docs/cub/Doxyfile b/docs/cub/Doxyfile index 6f2dec25329..f3fefa36a3a 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 From 26fbfe83b3599ffc7a3bb4d89f2de39b0844772c Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 8 Jul 2026 11:31:51 +0200 Subject: [PATCH 3/8] Order keyswords to fix Breathe bug --- cub/cub/device/dispatch/tuning/common.cuh | 4 ++-- .../tuning/tuning_adjacent_difference.cuh | 4 ++-- .../dispatch/tuning/tuning_batch_memcpy.cuh | 12 +++++----- .../dispatch/tuning/tuning_batched_topk.cuh | 16 ++++++------- .../device/dispatch/tuning/tuning_find.cuh | 4 ++-- .../tuning_find_bound_sorted_values.cuh | 4 ++-- cub/cub/device/dispatch/tuning/tuning_for.cuh | 4 ++-- .../dispatch/tuning/tuning_histogram.cuh | 4 ++-- .../device/dispatch/tuning/tuning_merge.cuh | 4 ++-- .../dispatch/tuning/tuning_merge_sort.cuh | 4 ++-- .../dispatch/tuning/tuning_radix_sort.cuh | 24 +++++++++---------- .../device/dispatch/tuning/tuning_reduce.cuh | 8 +++---- .../dispatch/tuning/tuning_reduce_by_key.cuh | 4 ++-- .../dispatch/tuning/tuning_rle_encode.cuh | 4 ++-- .../tuning/tuning_rle_non_trivial_runs.cuh | 4 ++-- .../device/dispatch/tuning/tuning_scan.cuh | 12 +++++----- .../dispatch/tuning/tuning_scan_by_key.cuh | 4 ++-- .../tuning/tuning_segmented_radix_sort.cuh | 4 ++-- .../tuning/tuning_segmented_reduce.cuh | 8 +++---- .../dispatch/tuning/tuning_segmented_scan.cuh | 8 +++---- .../dispatch/tuning/tuning_segmented_sort.cuh | 12 +++++----- .../dispatch/tuning/tuning_select_if.cuh | 8 +++---- .../tuning/tuning_three_way_partition.cuh | 4 ++-- .../device/dispatch/tuning/tuning_topk.cuh | 4 ++-- .../dispatch/tuning/tuning_transform.cuh | 20 ++++++++-------- .../dispatch/tuning/tuning_unique_by_key.cuh | 4 ++-- 26 files changed, 96 insertions(+), 96 deletions(-) diff --git a/cub/cub/device/dispatch/tuning/common.cuh b/cub/cub/device/dispatch/tuning/common.cuh index c7640a1abb9..894013f2b88 100644 --- a/cub/cub/device/dispatch/tuning/common.cuh +++ b/cub/cub/device/dispatch/tuning/common.cuh @@ -265,13 +265,13 @@ struct LookbackDelayPolicy 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 + [[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 constexpr friend bool + [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool operator!=(const LookbackDelayPolicy& lhs, const LookbackDelayPolicy& rhs) noexcept { return !(lhs == rhs); 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_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; // Date: Wed, 8 Jul 2026 11:38:36 +0200 Subject: [PATCH 4/8] Predefine _CCCL_HOSTED()=1 --- docs/cub/Doxyfile | 1 + docs/cudax/Doxyfile | 3 ++- docs/libcudacxx/Doxyfile | 3 ++- docs/thrust/Doxyfile | 3 ++- 4 files changed, 7 insertions(+), 3 deletions(-) diff --git a/docs/cub/Doxyfile b/docs/cub/Doxyfile index f3fefa36a3a..4d677c3cea4 100644 --- a/docs/cub/Doxyfile +++ b/docs/cub/Doxyfile @@ -170,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 From 40bf3752114c73f27ae68a7ccfa9131ceff5cb6e Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 8 Jul 2026 13:33:27 +0200 Subject: [PATCH 5/8] Fixes --- cub/cub/device/dispatch/tuning/common.cuh | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/cub/cub/device/dispatch/tuning/common.cuh b/cub/cub/device/dispatch/tuning/common.cuh index 894013f2b88..fbc1bc278c8 100644 --- a/cub/cub/device/dispatch/tuning/common.cuh +++ b/cub/cub/device/dispatch/tuning/common.cuh @@ -20,6 +20,7 @@ #include #include +#include #include #include #include @@ -203,7 +204,6 @@ enum class LookbackDelayAlgorithm __reduce_by_key //!< Internal }; -#if _CCCL_HOSTED() namespace detail { [[nodiscard]] constexpr const char* to_string(LookbackDelayAlgorithm algo) noexcept @@ -233,7 +233,6 @@ namespace detail } } } // namespace detail -#endif // _CCCL_HOSTED() #if _CCCL_HOSTED() inline ::std::ostream& operator<<(::std::ostream& os, LookbackDelayAlgorithm algo) From f38a9c82f2941567fbd41fc23f14ba6b60048da7 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 8 Jul 2026 17:30:16 +0200 Subject: [PATCH 6/8] Move default cases in to_string functions out of switches --- cub/cub/agent/agent_histogram.cuh | 3 +-- cub/cub/agent/agent_radix_sort_onesweep.cuh | 3 +-- cub/cub/block/block_load.cuh | 3 +-- cub/cub/block/block_radix_rank.cuh | 3 +-- cub/cub/block/block_reduce.cuh | 3 +-- cub/cub/block/block_scan.cuh | 3 +-- cub/cub/block/block_store.cuh | 3 +-- cub/cub/device/dispatch/tuning/common.cuh | 3 +-- cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh | 3 +-- cub/cub/device/dispatch/tuning/tuning_scan.cuh | 3 +-- cub/cub/device/dispatch/tuning/tuning_transform.cuh | 3 +-- cub/cub/thread/thread_load.cuh | 3 +-- cub/cub/warp/warp_load.cuh | 3 +-- cub/cub/warp/warp_store.cuh | 3 +-- 14 files changed, 14 insertions(+), 28 deletions(-) diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh index ab2abc62f15..93c3003ae86 100644 --- a/cub/cub/agent/agent_histogram.cuh +++ b/cub/cub/agent/agent_histogram.cuh @@ -53,9 +53,8 @@ namespace detail return "SMEM"; case BLEND: return "BLEND"; - default: - return ""; } + return ""; } } // namespace detail diff --git a/cub/cub/agent/agent_radix_sort_onesweep.cuh b/cub/cub/agent/agent_radix_sort_onesweep.cuh index 5980798d3db..29121089443 100644 --- a/cub/cub/agent/agent_radix_sort_onesweep.cuh +++ b/cub/cub/agent/agent_radix_sort_onesweep.cuh @@ -64,9 +64,8 @@ namespace detail return "RADIX_SORT_STORE_DIRECT"; case RADIX_SORT_STORE_ALIGNED: return "RADIX_SORT_STORE_ALIGNED"; - default: - return ""; } + return ""; } } // namespace detail #endif // _CCCL_HOSTED() diff --git a/cub/cub/block/block_load.cuh b/cub/cub/block/block_load.cuh index edd6d340f6a..cd377a8535f 100644 --- a/cub/cub/block/block_load.cuh +++ b/cub/cub/block/block_load.cuh @@ -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 ""; } + return ""; } } // namespace detail diff --git a/cub/cub/block/block_radix_rank.cuh b/cub/cub/block/block_radix_rank.cuh index 7482951c83d..3082031f99c 100644 --- a/cub/cub/block/block_radix_rank.cuh +++ b/cub/cub/block/block_radix_rank.cuh @@ -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 ""; } + return ""; } } // namespace detail diff --git a/cub/cub/block/block_reduce.cuh b/cub/cub/block/block_reduce.cuh index 6d7da26213b..9f2c22cda04 100644 --- a/cub/cub/block/block_reduce.cuh +++ b/cub/cub/block/block_reduce.cuh @@ -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 ""; } + return ""; } } // namespace detail diff --git a/cub/cub/block/block_scan.cuh b/cub/cub/block/block_scan.cuh index 195ea5658b9..333229a32ee 100644 --- a/cub/cub/block/block_scan.cuh +++ b/cub/cub/block/block_scan.cuh @@ -115,9 +115,8 @@ namespace detail return "BLOCK_SCAN_RAKING_MEMOIZE"; case BLOCK_SCAN_WARP_SCANS: return "BLOCK_SCAN_WARP_SCANS"; - default: - return ""; } + return ""; } } // namespace detail diff --git a/cub/cub/block/block_store.cuh b/cub/cub/block/block_store.cuh index c94dd0bac67..e6c1cee3795 100644 --- a/cub/cub/block/block_store.cuh +++ b/cub/cub/block/block_store.cuh @@ -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 ""; } + return ""; } } // namespace detail diff --git a/cub/cub/device/dispatch/tuning/common.cuh b/cub/cub/device/dispatch/tuning/common.cuh index fbc1bc278c8..964696eae7e 100644 --- a/cub/cub/device/dispatch/tuning/common.cuh +++ b/cub/cub/device/dispatch/tuning/common.cuh @@ -228,9 +228,8 @@ namespace detail return "LookbackDelayAlgorithm::exponential_backon"; case LookbackDelayAlgorithm::__reduce_by_key: return "LookbackDelayAlgorithm::__reduce_by_key"; - default: - return ""; } + return ""; } } // namespace detail diff --git a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh index 06daa1adbd2..e242689569e 100644 --- a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh @@ -49,9 +49,8 @@ namespace detail return "RadixSortAlgorithm::multi_pass"; case RadixSortAlgorithm::onesweep: return "RadixSortAlgorithm::onesweep"; - default: - return ""; } + return ""; } } // namespace detail #endif // _CCCL_HOSTED() diff --git a/cub/cub/device/dispatch/tuning/tuning_scan.cuh b/cub/cub/device/dispatch/tuning/tuning_scan.cuh index 427d2e05c53..7a9aefab364 100644 --- a/cub/cub/device/dispatch/tuning/tuning_scan.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_scan.cuh @@ -63,9 +63,8 @@ namespace detail return "ScanAlgorithm::lookback"; case ScanAlgorithm::lookahead: return "ScanAlgorithm::lookahead"; - default: - return ""; } + return ""; } } // namespace detail #endif // _CCCL_HOSTED() diff --git a/cub/cub/device/dispatch/tuning/tuning_transform.cuh b/cub/cub/device/dispatch/tuning/tuning_transform.cuh index 76af2a2ebb4..5b1fcd002db 100644 --- a/cub/cub/device/dispatch/tuning/tuning_transform.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_transform.cuh @@ -59,9 +59,8 @@ namespace detail return "TransformAlgorithm::ldgsts"; case TransformAlgorithm::ublkcp: return "TransformAlgorithm::ublkcp"; - default: - return ""; } + return ""; } } // namespace detail #endif // _CCCL_HOSTED() diff --git a/cub/cub/thread/thread_load.cuh b/cub/cub/thread/thread_load.cuh index a881322e1b2..8af1193d3ce 100644 --- a/cub/cub/thread/thread_load.cuh +++ b/cub/cub/thread/thread_load.cuh @@ -69,9 +69,8 @@ namespace detail return "LOAD_LDG"; case LOAD_VOLATILE: return "LOAD_VOLATILE"; - default: - return ""; } + return ""; } } // namespace detail #endif // _CCCL_HOSTED() diff --git a/cub/cub/warp/warp_load.cuh b/cub/cub/warp/warp_load.cuh index 3ecdc9d5e2f..dec85fd70ec 100644 --- a/cub/cub/warp/warp_load.cuh +++ b/cub/cub/warp/warp_load.cuh @@ -124,9 +124,8 @@ namespace detail return "WARP_LOAD_VECTORIZE"; case WARP_LOAD_TRANSPOSE: return "WARP_LOAD_TRANSPOSE"; - default: - return ""; } + return ""; } } // namespace detail #endif // _CCCL_HOSTED() diff --git a/cub/cub/warp/warp_store.cuh b/cub/cub/warp/warp_store.cuh index 308e9c6aaef..570dd0f05ec 100644 --- a/cub/cub/warp/warp_store.cuh +++ b/cub/cub/warp/warp_store.cuh @@ -128,9 +128,8 @@ namespace detail return "WARP_STORE_VECTORIZE"; case WARP_STORE_TRANSPOSE: return "WARP_STORE_TRANSPOSE"; - default: - return ""; } + return ""; } } // namespace detail From 07d5e75107859835b09574926405e6176800766f Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 8 Jul 2026 17:34:06 +0200 Subject: [PATCH 7/8] Fully qualify std::formatter specializations --- cub/cub/agent/agent_histogram.cuh | 2 +- cub/cub/agent/agent_radix_sort_onesweep.cuh | 2 +- cub/cub/block/block_load.cuh | 2 +- cub/cub/block/block_radix_rank.cuh | 2 +- cub/cub/block/block_reduce.cuh | 2 +- cub/cub/block/block_scan.cuh | 2 +- cub/cub/block/block_store.cuh | 2 +- cub/cub/device/dispatch/tuning/common.cuh | 2 +- cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh | 2 +- cub/cub/device/dispatch/tuning/tuning_scan.cuh | 2 +- cub/cub/device/dispatch/tuning/tuning_transform.cuh | 2 +- cub/cub/thread/thread_load.cuh | 2 +- cub/cub/warp/warp_load.cuh | 2 +- cub/cub/warp/warp_store.cuh | 2 +- 14 files changed, 14 insertions(+), 14 deletions(-) diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh index 93c3003ae86..4f43eb52996 100644 --- a/cub/cub/agent/agent_histogram.cuh +++ b/cub/cub/agent/agent_histogram.cuh @@ -68,7 +68,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template 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 29121089443..01edd65e520 100644 --- a/cub/cub/agent/agent_radix_sort_onesweep.cuh +++ b/cub/cub/agent/agent_radix_sort_onesweep.cuh @@ -81,7 +81,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template 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 cd377a8535f..83d9f89ea64 100644 --- a/cub/cub/block/block_load.cuh +++ b/cub/cub/block/block_load.cuh @@ -763,7 +763,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template 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 3082031f99c..9781f8606ce 100644 --- a/cub/cub/block/block_radix_rank.cuh +++ b/cub/cub/block/block_radix_rank.cuh @@ -105,7 +105,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template 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 9f2c22cda04..b1814f8d77a 100644 --- a/cub/cub/block/block_reduce.cuh +++ b/cub/cub/block/block_reduce.cuh @@ -181,7 +181,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template 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 333229a32ee..309f52b0740 100644 --- a/cub/cub/block/block_scan.cuh +++ b/cub/cub/block/block_scan.cuh @@ -130,7 +130,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template 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 e6c1cee3795..889317fd861 100644 --- a/cub/cub/block/block_store.cuh +++ b/cub/cub/block/block_store.cuh @@ -579,7 +579,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template auto format(const CUB_NS_QUALIFIER::BlockStoreAlgorithm& algo, FmtCtx& ctx) const diff --git a/cub/cub/device/dispatch/tuning/common.cuh b/cub/cub/device/dispatch/tuning/common.cuh index 964696eae7e..921e4d53355 100644 --- a/cub/cub/device/dispatch/tuning/common.cuh +++ b/cub/cub/device/dispatch/tuning/common.cuh @@ -244,7 +244,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template auto format(const CUB_NS_QUALIFIER::LookbackDelayAlgorithm& algo, FmtCtx& ctx) const diff --git a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh index e242689569e..0da8190467b 100644 --- a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh @@ -66,7 +66,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template auto format(const CUB_NS_QUALIFIER::RadixSortAlgorithm& algo, FmtCtx& ctx) const diff --git a/cub/cub/device/dispatch/tuning/tuning_scan.cuh b/cub/cub/device/dispatch/tuning/tuning_scan.cuh index 7a9aefab364..ff5a9ff4954 100644 --- a/cub/cub/device/dispatch/tuning/tuning_scan.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_scan.cuh @@ -80,7 +80,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template auto format(const CUB_NS_QUALIFIER::ScanAlgorithm& algo, FmtCtx& ctx) const diff --git a/cub/cub/device/dispatch/tuning/tuning_transform.cuh b/cub/cub/device/dispatch/tuning/tuning_transform.cuh index 5b1fcd002db..bdd0ac2e367 100644 --- a/cub/cub/device/dispatch/tuning/tuning_transform.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_transform.cuh @@ -76,7 +76,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template auto format(const CUB_NS_QUALIFIER::TransformAlgorithm& algo, FmtCtx& ctx) const diff --git a/cub/cub/thread/thread_load.cuh b/cub/cub/thread/thread_load.cuh index 8af1193d3ce..9607ff42cb6 100644 --- a/cub/cub/thread/thread_load.cuh +++ b/cub/cub/thread/thread_load.cuh @@ -86,7 +86,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template 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 dec85fd70ec..93ed880306a 100644 --- a/cub/cub/warp/warp_load.cuh +++ b/cub/cub/warp/warp_load.cuh @@ -141,7 +141,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template 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 570dd0f05ec..35c603cf719 100644 --- a/cub/cub/warp/warp_store.cuh +++ b/cub/cub/warp/warp_store.cuh @@ -143,7 +143,7 @@ CUB_NAMESPACE_END #if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) template <::cuda::std::same_as CharT> -struct std::formatter : formatter +struct ::std::formatter : formatter { template auto format(const CUB_NS_QUALIFIER::WarpStoreAlgorithm& algo, FmtCtx& ctx) const From 704b9418f6900605274cda75c667d622b22a05d1 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 8 Jul 2026 17:36:38 +0200 Subject: [PATCH 8/8] _CCCL_API --- cub/cub/agent/agent_histogram.cuh | 2 +- cub/cub/agent/agent_radix_sort_onesweep.cuh | 2 +- cub/cub/block/block_load.cuh | 2 +- cub/cub/block/block_radix_rank.cuh | 2 +- cub/cub/block/block_reduce.cuh | 2 +- cub/cub/block/block_scan.cuh | 2 +- cub/cub/block/block_store.cuh | 2 +- cub/cub/device/dispatch/tuning/common.cuh | 2 +- cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh | 2 +- cub/cub/device/dispatch/tuning/tuning_scan.cuh | 2 +- cub/cub/device/dispatch/tuning/tuning_transform.cuh | 2 +- cub/cub/thread/thread_load.cuh | 2 +- cub/cub/warp/warp_load.cuh | 2 +- cub/cub/warp/warp_store.cuh | 2 +- 14 files changed, 14 insertions(+), 14 deletions(-) diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh index 4f43eb52996..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) { diff --git a/cub/cub/agent/agent_radix_sort_onesweep.cuh b/cub/cub/agent/agent_radix_sort_onesweep.cuh index 01edd65e520..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) { diff --git a/cub/cub/block/block_load.cuh b/cub/cub/block/block_load.cuh index 83d9f89ea64..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) { diff --git a/cub/cub/block/block_radix_rank.cuh b/cub/cub/block/block_radix_rank.cuh index 9781f8606ce..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) { diff --git a/cub/cub/block/block_reduce.cuh b/cub/cub/block/block_reduce.cuh index b1814f8d77a..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) { diff --git a/cub/cub/block/block_scan.cuh b/cub/cub/block/block_scan.cuh index 309f52b0740..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) { diff --git a/cub/cub/block/block_store.cuh b/cub/cub/block/block_store.cuh index 889317fd861..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) { diff --git a/cub/cub/device/dispatch/tuning/common.cuh b/cub/cub/device/dispatch/tuning/common.cuh index 921e4d53355..02e524ef464 100644 --- a/cub/cub/device/dispatch/tuning/common.cuh +++ b/cub/cub/device/dispatch/tuning/common.cuh @@ -206,7 +206,7 @@ enum class LookbackDelayAlgorithm namespace detail { -[[nodiscard]] constexpr const char* to_string(LookbackDelayAlgorithm algo) noexcept +_CCCL_API [[nodiscard]] constexpr const char* to_string(LookbackDelayAlgorithm algo) noexcept { switch (algo) { diff --git a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh index 0da8190467b..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) { diff --git a/cub/cub/device/dispatch/tuning/tuning_scan.cuh b/cub/cub/device/dispatch/tuning/tuning_scan.cuh index ff5a9ff4954..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) { diff --git a/cub/cub/device/dispatch/tuning/tuning_transform.cuh b/cub/cub/device/dispatch/tuning/tuning_transform.cuh index bdd0ac2e367..7ce6d997a53 100644 --- a/cub/cub/device/dispatch/tuning/tuning_transform.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_transform.cuh @@ -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) { diff --git a/cub/cub/thread/thread_load.cuh b/cub/cub/thread/thread_load.cuh index 9607ff42cb6..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) { diff --git a/cub/cub/warp/warp_load.cuh b/cub/cub/warp/warp_load.cuh index 93ed880306a..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) { diff --git a/cub/cub/warp/warp_store.cuh b/cub/cub/warp/warp_store.cuh index 35c603cf719..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) {