diff --git a/cub/cub/agent/agent_batched_topk.cuh b/cub/cub/agent/agent_batched_topk.cuh index 0bdaa8f38ef..c5ec20d26cd 100644 --- a/cub/cub/agent/agent_batched_topk.cuh +++ b/cub/cub/agent/agent_batched_topk.cuh @@ -95,7 +95,7 @@ struct agent_batched_topk_worker_per_segment multi_worker_per_segment_policy.threads_per_block * multi_worker_per_segment_policy.items_per_thread; // Check if there could be large segments present - static constexpr bool only_small_segments = ::cuda::__argument::__traits::max <= tile_size; + static constexpr bool only_small_segments = ::cuda::__argument::__traits::highest <= tile_size; // Check if we are dealing with keys-only or key-value pairs static constexpr bool is_keys_only = ::cuda::std::is_same_v; @@ -224,7 +224,7 @@ struct agent_batched_topk_worker_per_segment const key_t padding_key = (direction == detail::topk::select::max) ? ::cuda::std::numeric_limits::lowest() - : ::cuda::std::numeric_limits::max(); + : (::cuda::std::numeric_limits::max)(); // Dereference iterator-of-iterators to get the segment specific iterator auto block_keys_in = d_key_segments_it[segment_id]; diff --git a/cub/cub/device/dispatch/dispatch_batched_topk.cuh b/cub/cub/device/dispatch/dispatch_batched_topk.cuh index 3d8293742eb..d0f2d4eed0a 100644 --- a/cub/cub/device/dispatch/dispatch_batched_topk.cuh +++ b/cub/cub/device/dispatch/dispatch_batched_topk.cuh @@ -126,7 +126,7 @@ template >, it_value_t>, ::cuda::std::int64_t, - ::cuda::__argument::__traits::max>> + ::cuda::__argument::__traits::highest>> #if _CCCL_HAS_CONCEPTS() requires batched_topk_policy_selector #endif // _CCCL_HAS_CONCEPTS() @@ -173,7 +173,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( static constexpr bool any_small_segments = ::cuda::__argument::__traits::lowest <= worker_per_segment_tile_size; static constexpr bool only_small_segments = - ::cuda::__argument::__traits::max <= worker_per_segment_tile_size; + ::cuda::__argument::__traits::highest <= worker_per_segment_tile_size; // Allocation layout: // only_small_segments: [0] dummy. @@ -341,7 +341,7 @@ template >, it_value_t>, ::cuda::std::int64_t, - ::cuda::__argument::__traits::max>; + ::cuda::__argument::__traits::highest>; return detail::dispatch_with_env_and_tuning( env, [&](auto policy_selector, void* d_temp_storage, size_t& temp_storage_bytes, cudaStream_t stream) { return dispatch( diff --git a/cub/cub/device/dispatch/kernels/kernel_batched_topk.cuh b/cub/cub/device/dispatch/kernels/kernel_batched_topk.cuh index 63ad24aaeac..1ff50dfaf67 100644 --- a/cub/cub/device/dispatch/kernels/kernel_batched_topk.cuh +++ b/cub/cub/device/dispatch/kernels/kernel_batched_topk.cuh @@ -39,7 +39,7 @@ private: worker_policy worker_per_segment_policy; multi_worker_policy multi_worker_per_segment_policy; }; - static constexpr ::cuda::std::int64_t max_segment_size = ::cuda::__argument::__traits::max; + static constexpr ::cuda::std::int64_t max_segment_size = ::cuda::__argument::__traits::highest; static constexpr batched_topk_policy active_policy = current_policy(); template @@ -151,7 +151,7 @@ __launch_bounds__(int( LargeSegmentTileOffsetT>::agent_t; // Static Assertions (Constraints) - static_assert(agent_t::tile_size >= ::cuda::__argument::__traits::max, + static_assert(agent_t::tile_size >= ::cuda::__argument::__traits::highest, "Block size exceeds maximum segment size supported by SegmentSizeParameterT"); static_assert(sizeof(typename agent_t::TempStorage) <= max_smem_per_block, "Static shared memory per block must not exceed 48KB limit."); diff --git a/libcudacxx/include/cuda/__argument/argument.h b/libcudacxx/include/cuda/__argument/argument.h index bd634fd4943..fb514d22648 100644 --- a/libcudacxx/include/cuda/__argument/argument.h +++ b/libcudacxx/include/cuda/__argument/argument.h @@ -173,9 +173,9 @@ _CCCL_API constexpr bool __static_bound_in_range() noexcept template inline constexpr bool __valid_static_bounds_v = true; -template -inline constexpr bool __valid_static_bounds_v<_ElementType, __static_bounds<_Lowest, _Max>> = - __static_bound_in_range<_ElementType, _Lowest>() && __static_bound_in_range<_ElementType, _Max>(); +template +inline constexpr bool __valid_static_bounds_v<_ElementType, __static_bounds<_Lowest, _Highest>> = + __static_bound_in_range<_ElementType, _Lowest>() && __static_bound_in_range<_ElementType, _Highest>(); template _CCCL_API constexpr _ElementType __wrapper_static_lowest() noexcept @@ -191,11 +191,11 @@ _CCCL_API constexpr _ElementType __wrapper_static_lowest() noexcept } template -_CCCL_API constexpr _ElementType __wrapper_static_max() noexcept +_CCCL_API constexpr _ElementType __wrapper_static_highest() noexcept { if constexpr (::cuda::std::is_same_v<_StaticBounds, __no_bounds>) { - return ::cuda::std::numeric_limits<_ElementType>::max(); + return (::cuda::std::numeric_limits<_ElementType>::max)(); } else { @@ -211,17 +211,17 @@ _CCCL_API constexpr _ElementType __effective_lowest(__runtime_bounds<_ElementTyp } template -_CCCL_API constexpr _ElementType __effective_max(__runtime_bounds<_ElementType> __runtime_bounds) noexcept +_CCCL_API constexpr _ElementType __effective_highest(__runtime_bounds<_ElementType> __runtime_bounds) noexcept { - auto __static_max = __wrapper_static_max<_ElementType, _StaticBounds>(); - return __static_max < __runtime_bounds.upper() ? __static_max : __runtime_bounds.upper(); + auto __static_highest = __wrapper_static_highest<_ElementType, _StaticBounds>(); + return __static_highest < __runtime_bounds.upper() ? __static_highest : __runtime_bounds.upper(); } template _CCCL_API constexpr bool __has_bounds_intersection(__runtime_bounds<_ElementType> __runtime_bounds) noexcept { return __effective_lowest<_ElementType, _StaticBounds>(__runtime_bounds) - <= __effective_max<_ElementType, _StaticBounds>(__runtime_bounds); + <= __effective_highest<_ElementType, _StaticBounds>(__runtime_bounds); } template @@ -240,8 +240,8 @@ _CCCL_API constexpr void __validate_static_element_bounds([[maybe_unused]] const { _CCCL_ASSERT((__val >= __wrapper_static_lowest<_ElementType, _StaticBounds>()), "immediate argument value is below static lowest bound"); - _CCCL_ASSERT((__val <= __wrapper_static_max<_ElementType, _StaticBounds>()), - "immediate argument value is above static max bound"); + _CCCL_ASSERT((__val <= __wrapper_static_highest<_ElementType, _StaticBounds>()), + "immediate argument value is above static highest bound"); } } @@ -295,9 +295,9 @@ struct __immediate }; #ifndef _CCCL_DOXYGEN_INVOKED -template -_CCCL_HOST_DEVICE __immediate(_Arg, __static_bounds<_Lowest, _Max>) - -> __immediate<_Arg, __static_bounds<_Lowest, _Max>>; +template +_CCCL_HOST_DEVICE __immediate(_Arg, __static_bounds<_Lowest, _Highest>) + -> __immediate<_Arg, __static_bounds<_Lowest, _Highest>>; #endif // _CCCL_DOXYGEN_INVOKED @@ -383,17 +383,17 @@ struct __immediate_sequence }; #ifndef _CCCL_DOXYGEN_INVOKED -template -_CCCL_HOST_DEVICE __immediate_sequence(_Arg, __static_bounds<_Lowest, _Max>) - -> __immediate_sequence<_Arg, __static_bounds<_Lowest, _Max>>; +template +_CCCL_HOST_DEVICE __immediate_sequence(_Arg, __static_bounds<_Lowest, _Highest>) + -> __immediate_sequence<_Arg, __static_bounds<_Lowest, _Highest>>; -template -_CCCL_HOST_DEVICE __immediate_sequence(_Arg, __static_bounds<_Lowest, _Max>, __runtime_bounds<_Tp>) - -> __immediate_sequence<_Arg, __static_bounds<_Lowest, _Max>>; +template +_CCCL_HOST_DEVICE __immediate_sequence(_Arg, __static_bounds<_Lowest, _Highest>, __runtime_bounds<_Tp>) + -> __immediate_sequence<_Arg, __static_bounds<_Lowest, _Highest>>; -template -_CCCL_HOST_DEVICE __immediate_sequence(_Arg, __runtime_bounds<_Tp>, __static_bounds<_Lowest, _Max>) - -> __immediate_sequence<_Arg, __static_bounds<_Lowest, _Max>>; +template +_CCCL_HOST_DEVICE __immediate_sequence(_Arg, __runtime_bounds<_Tp>, __static_bounds<_Lowest, _Highest>) + -> __immediate_sequence<_Arg, __static_bounds<_Lowest, _Highest>>; #endif // _CCCL_DOXYGEN_INVOKED // ===================================================================== @@ -460,19 +460,20 @@ struct __deferred : __deferred_base<_Arg, _StaticBounds> template _CCCL_HOST_DEVICE __deferred(_Arg) -> __deferred<_Arg>; -template -_CCCL_HOST_DEVICE __deferred(_Arg, __static_bounds<_Lowest, _Max>) -> __deferred<_Arg, __static_bounds<_Lowest, _Max>>; +template +_CCCL_HOST_DEVICE __deferred(_Arg, __static_bounds<_Lowest, _Highest>) + -> __deferred<_Arg, __static_bounds<_Lowest, _Highest>>; template _CCCL_HOST_DEVICE __deferred(_Arg, __runtime_bounds<_Tp>) -> __deferred<_Arg>; -template -_CCCL_HOST_DEVICE __deferred(_Arg, __static_bounds<_Lowest, _Max>, __runtime_bounds<_Tp>) - -> __deferred<_Arg, __static_bounds<_Lowest, _Max>>; +template +_CCCL_HOST_DEVICE __deferred(_Arg, __static_bounds<_Lowest, _Highest>, __runtime_bounds<_Tp>) + -> __deferred<_Arg, __static_bounds<_Lowest, _Highest>>; -template -_CCCL_HOST_DEVICE __deferred(_Arg, __runtime_bounds<_Tp>, __static_bounds<_Lowest, _Max>) - -> __deferred<_Arg, __static_bounds<_Lowest, _Max>>; +template +_CCCL_HOST_DEVICE __deferred(_Arg, __runtime_bounds<_Tp>, __static_bounds<_Lowest, _Highest>) + -> __deferred<_Arg, __static_bounds<_Lowest, _Highest>>; #endif // _CCCL_DOXYGEN_INVOKED //! @brief Wraps a reference to a sequence of values that is potentially not available at API call time but will be @@ -487,20 +488,20 @@ struct __deferred_sequence : __deferred_base<_Arg, _StaticBounds> template _CCCL_HOST_DEVICE __deferred_sequence(_Arg) -> __deferred_sequence<_Arg>; -template -_CCCL_HOST_DEVICE __deferred_sequence(_Arg, __static_bounds<_Lowest, _Max>) - -> __deferred_sequence<_Arg, __static_bounds<_Lowest, _Max>>; +template +_CCCL_HOST_DEVICE __deferred_sequence(_Arg, __static_bounds<_Lowest, _Highest>) + -> __deferred_sequence<_Arg, __static_bounds<_Lowest, _Highest>>; template _CCCL_HOST_DEVICE __deferred_sequence(_Arg, __runtime_bounds<_Tp>) -> __deferred_sequence<_Arg>; -template -_CCCL_HOST_DEVICE __deferred_sequence(_Arg, __static_bounds<_Lowest, _Max>, __runtime_bounds<_Tp>) - -> __deferred_sequence<_Arg, __static_bounds<_Lowest, _Max>>; +template +_CCCL_HOST_DEVICE __deferred_sequence(_Arg, __static_bounds<_Lowest, _Highest>, __runtime_bounds<_Tp>) + -> __deferred_sequence<_Arg, __static_bounds<_Lowest, _Highest>>; -template -_CCCL_HOST_DEVICE __deferred_sequence(_Arg, __runtime_bounds<_Tp>, __static_bounds<_Lowest, _Max>) - -> __deferred_sequence<_Arg, __static_bounds<_Lowest, _Max>>; +template +_CCCL_HOST_DEVICE __deferred_sequence(_Arg, __runtime_bounds<_Tp>, __static_bounds<_Lowest, _Highest>) + -> __deferred_sequence<_Arg, __static_bounds<_Lowest, _Highest>>; #endif // _CCCL_DOXYGEN_INVOKED // ===================================================================== @@ -622,7 +623,7 @@ _CCCL_API constexpr auto __constant_compute_lowest() noexcept } template -_CCCL_API constexpr auto __constant_compute_max() noexcept +_CCCL_API constexpr auto __constant_compute_highest() noexcept { return _Value; } @@ -642,7 +643,7 @@ _CCCL_API constexpr auto __constant_sequence_compute_lowest() noexcept } template -_CCCL_API constexpr auto __constant_sequence_compute_max() noexcept +_CCCL_API constexpr auto __constant_sequence_compute_highest() noexcept { using _ElementType = __element_type_of_t<::cuda::std::remove_cvref_t>; auto __first = _Value.begin(); @@ -650,7 +651,7 @@ _CCCL_API constexpr auto __constant_sequence_compute_max() noexcept if (__first == __last) { - return ::cuda::std::numeric_limits<_ElementType>::max(); + return (::cuda::std::numeric_limits<_ElementType>::max)(); } return static_cast<_ElementType>(*::cuda::std::max_element(__first, __last)); } @@ -661,7 +662,7 @@ _CCCL_API constexpr auto __constant_sequence_compute_max() noexcept //! @brief Traits for argument wrappers and plain argument values. //! -//! Models @c numeric_limits for bounds: @c lowest is the lower bound, @c max is the upper bound. +//! Models @c numeric_limits for bounds: @c lowest is the lower bound, @c highest is the upper bound. //! Use in @c if @c constexpr for compile-time dispatch based on bounds. template struct __traits_impl @@ -672,7 +673,7 @@ struct __traits_impl static constexpr bool is_deferred = false; static constexpr bool is_single_value = __is_single_value_v<_Tp>; static constexpr element_type lowest = ::cuda::std::numeric_limits::lowest(); - static constexpr element_type max = ::cuda::std::numeric_limits::max(); + static constexpr element_type highest = (::cuda::std::numeric_limits::max)(); }; template @@ -684,7 +685,7 @@ struct __traits_impl<__constant<_Value>> static constexpr bool is_deferred = false; static constexpr bool is_single_value = true; static constexpr element_type lowest = __constant_compute_lowest<_Value>(); - static constexpr element_type max = __constant_compute_max<_Value>(); + static constexpr element_type highest = __constant_compute_highest<_Value>(); }; template @@ -699,7 +700,7 @@ struct __traits_impl<__immediate<_Arg, _StaticBounds>> static constexpr bool is_deferred = false; static constexpr bool is_single_value = true; static constexpr element_type lowest = __wrapper_static_lowest(); - static constexpr element_type max = __wrapper_static_max(); + static constexpr element_type highest = __wrapper_static_highest(); }; template @@ -712,7 +713,7 @@ struct __traits_impl<__constant_sequence<_Value>> static constexpr bool is_deferred = false; static constexpr bool is_single_value = false; static constexpr element_type lowest = __constant_sequence_compute_lowest<_Value>(); - static constexpr element_type max = __constant_sequence_compute_max<_Value>(); + static constexpr element_type highest = __constant_sequence_compute_highest<_Value>(); }; template @@ -728,7 +729,7 @@ struct __traits_impl<__immediate_sequence<_Arg, _StaticBounds>> static constexpr bool is_deferred = false; static constexpr bool is_single_value = false; static constexpr element_type lowest = __wrapper_static_lowest(); - static constexpr element_type max = __wrapper_static_max(); + static constexpr element_type highest = __wrapper_static_highest(); }; template @@ -743,7 +744,7 @@ struct __traits_impl<__deferred<_Arg, _StaticBounds>> static constexpr bool is_deferred = true; static constexpr bool is_single_value = true; static constexpr element_type lowest = __wrapper_static_lowest(); - static constexpr element_type max = __wrapper_static_max(); + static constexpr element_type highest = __wrapper_static_highest(); }; template @@ -758,7 +759,7 @@ struct __traits_impl<__deferred_sequence<_Arg, _StaticBounds>> static constexpr bool is_deferred = true; static constexpr bool is_single_value = false; static constexpr element_type lowest = __wrapper_static_lowest(); - static constexpr element_type max = __wrapper_static_max(); + static constexpr element_type highest = __wrapper_static_highest(); }; template @@ -766,7 +767,7 @@ struct __traits : __traits_impl<::cuda::std::remove_cvref_t<_Tp>> {}; // ===================================================================== -// __lowest_ / __max_ — free functions +// __lowest_ / __highest_ — free functions // ===================================================================== //! @brief Returns the effective lowest bound, combining static and runtime bounds. @@ -819,54 +820,54 @@ template return __effective_lowest<_ET, _StaticBounds>(__arg.__runtime_bounds_); } -//! @brief Returns the effective max bound, combining static and runtime bounds. +//! @brief Returns the effective highest bound, combining static and runtime bounds. _CCCL_TEMPLATE(class _Tp) _CCCL_REQUIRES((!__is_wrapper_v<::cuda::std::remove_cv_t<_Tp>>) ) -[[nodiscard]] _CCCL_API constexpr auto __max_(_Tp) noexcept +[[nodiscard]] _CCCL_API constexpr auto __highest_(_Tp) noexcept { - return ::cuda::std::numeric_limits<__element_type_of_t<_Tp>>::max(); + return (::cuda::std::numeric_limits<__element_type_of_t<_Tp>>::max)(); } template -[[nodiscard]] _CCCL_API constexpr auto __max_(__constant<_Value>) noexcept +[[nodiscard]] _CCCL_API constexpr auto __highest_(__constant<_Value>) noexcept { - return __constant_compute_max<_Value>(); + return __constant_compute_highest<_Value>(); } template -[[nodiscard]] _CCCL_API constexpr auto __max_(__constant_sequence<_Value>) noexcept +[[nodiscard]] _CCCL_API constexpr auto __highest_(__constant_sequence<_Value>) noexcept { - return __constant_sequence_compute_max<_Value>(); + return __constant_sequence_compute_highest<_Value>(); } template -[[nodiscard]] _CCCL_API constexpr auto __max_(__immediate<_Arg, _StaticBounds> __arg) noexcept +[[nodiscard]] _CCCL_API constexpr auto __highest_(__immediate<_Arg, _StaticBounds> __arg) noexcept { return __arg.__arg_; } template -[[nodiscard]] _CCCL_API constexpr auto __max_(__immediate_sequence<_Arg, _StaticBounds> __arg) noexcept +[[nodiscard]] _CCCL_API constexpr auto __highest_(__immediate_sequence<_Arg, _StaticBounds> __arg) noexcept { using _ET = __element_type_of_t<_Arg>; __validate_bounds_intersection<_ET, _StaticBounds>(__arg.__runtime_bounds_); - return __effective_max<_ET, _StaticBounds>(__arg.__runtime_bounds_); + return __effective_highest<_ET, _StaticBounds>(__arg.__runtime_bounds_); } template -[[nodiscard]] _CCCL_API constexpr auto __max_(__deferred<_Arg, _StaticBounds> __arg) noexcept +[[nodiscard]] _CCCL_API constexpr auto __highest_(__deferred<_Arg, _StaticBounds> __arg) noexcept { using _ET = __element_type_of_t<_Arg>; __validate_bounds_intersection<_ET, _StaticBounds>(__arg.__runtime_bounds_); - return __effective_max<_ET, _StaticBounds>(__arg.__runtime_bounds_); + return __effective_highest<_ET, _StaticBounds>(__arg.__runtime_bounds_); } template -[[nodiscard]] _CCCL_API constexpr auto __max_(__deferred_sequence<_Arg, _StaticBounds> __arg) noexcept +[[nodiscard]] _CCCL_API constexpr auto __highest_(__deferred_sequence<_Arg, _StaticBounds> __arg) noexcept { using _ET = __element_type_of_t<_Arg>; __validate_bounds_intersection<_ET, _StaticBounds>(__arg.__runtime_bounds_); - return __effective_max<_ET, _StaticBounds>(__arg.__runtime_bounds_); + return __effective_highest<_ET, _StaticBounds>(__arg.__runtime_bounds_); } _CCCL_END_NAMESPACE_CUDA_ARGUMENT diff --git a/libcudacxx/include/cuda/__argument/argument_bounds.h b/libcudacxx/include/cuda/__argument/argument_bounds.h index b26d1a383a2..11c46c74417 100644 --- a/libcudacxx/include/cuda/__argument/argument_bounds.h +++ b/libcudacxx/include/cuda/__argument/argument_bounds.h @@ -77,7 +77,7 @@ template struct __runtime_bounds { _Tp __lower_ = ::cuda::std::numeric_limits<_Tp>::lowest(); - _Tp __upper_ = ::cuda::std::numeric_limits<_Tp>::max(); + _Tp __upper_ = (::cuda::std::numeric_limits<_Tp>::max)(); constexpr __runtime_bounds() noexcept = default; diff --git a/libcudacxx/test/libcudacxx/cuda/argument/argument_traits.pass.cpp b/libcudacxx/test/libcudacxx/cuda/argument/argument_traits.pass.cpp index 3c448dfb1d1..cdcc6665747 100644 --- a/libcudacxx/test/libcudacxx/cuda/argument/argument_traits.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/argument/argument_traits.pass.cpp @@ -137,36 +137,37 @@ TEST_FUNC void test() cuda::std::array>); #endif // TEST_HAS_CLASS_NTTP - // --- argument_traits: lowest / max --- + // --- argument_traits: lowest / highest --- static_assert(cuda::__argument::__traits::lowest == cuda::std::numeric_limits::lowest()); - static_assert(cuda::__argument::__traits::max == cuda::std::numeric_limits::max()); + static_assert(cuda::__argument::__traits::highest == (cuda::std::numeric_limits::max)()); static_assert(cuda::__argument::__traits::lowest == cuda::std::numeric_limits::lowest()); - static_assert(cuda::__argument::__traits::max == cuda::std::numeric_limits::max()); + static_assert(cuda::__argument::__traits::highest == (cuda::std::numeric_limits::max)()); static_assert(cuda::__argument::__traits::lowest == cuda::std::numeric_limits::lowest()); - static_assert(cuda::__argument::__traits::max == cuda::std::numeric_limits::max()); + static_assert(cuda::__argument::__traits::highest == (cuda::std::numeric_limits::max)()); static_assert( cuda::__argument::__traits>>::lowest == 1); static_assert( - cuda::__argument::__traits>&>::max == 8); + cuda::__argument::__traits>&>::highest + == 8); static_assert( cuda::__argument::__traits< - cuda::__argument::__immediate_sequence, cuda::__argument::__static_bounds<1, 8>>>::max + cuda::__argument::__immediate_sequence, cuda::__argument::__static_bounds<1, 8>>>::highest == 8); #if TEST_HAS_CLASS_NTTP static_assert( cuda::__argument::__traits{3, 1, 2}>>::lowest == 1); static_assert( - cuda::__argument::__traits{3, 1, 2}>>::max == 3); + cuda::__argument::__traits{3, 1, 2}>>::highest == 3); #endif // TEST_HAS_CLASS_NTTP // --- Free function bounds on plain values --- static_assert(cuda::__argument::__lowest_(42) == cuda::std::numeric_limits::lowest()); - static_assert(cuda::__argument::__max_(42) == cuda::std::numeric_limits::max()); + static_assert(cuda::__argument::__highest_(42) == (cuda::std::numeric_limits::max)()); static_assert(cuda::__argument::__lowest_(1.0f) == cuda::std::numeric_limits::lowest()); - static_assert(cuda::__argument::__max_(1.0f) == cuda::std::numeric_limits::max()); + static_assert(cuda::__argument::__highest_(1.0f) == (cuda::std::numeric_limits::max)()); // --- Scalar and sequence wrappers expose distinct single-value traits --- diff --git a/libcudacxx/test/libcudacxx/cuda/argument/deferred_argument.pass.cpp b/libcudacxx/test/libcudacxx/cuda/argument/deferred_argument.pass.cpp index ca5304e140f..9949e0013b4 100644 --- a/libcudacxx/test/libcudacxx/cuda/argument/deferred_argument.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/argument/deferred_argument.pass.cpp @@ -25,7 +25,7 @@ TEST_FUNC constexpr bool test() auto def = cuda::__argument::__deferred{cuda::std::span{&val, 1}}; assert(cuda::__argument::__unwrap(def)[0] == 42); static_assert(cuda::__argument::__traits::lowest == cuda::std::numeric_limits::lowest()); - static_assert(cuda::__argument::__traits::max == cuda::std::numeric_limits::max()); + static_assert(cuda::__argument::__traits::highest == (cuda::std::numeric_limits::max)()); } // Deferred single value with static bounds @@ -34,7 +34,7 @@ TEST_FUNC constexpr bool test() auto def = cuda::__argument::__deferred{cuda::std::span{&val, 1}, cuda::__argument::__bounds<1, 1000>()}; assert(cuda::__argument::__unwrap(def)[0] == 42); static_assert(cuda::__argument::__traits::lowest == 1); - static_assert(cuda::__argument::__traits::max == 1000); + static_assert(cuda::__argument::__traits::highest == 1000); } // Deferred single value via pointer @@ -42,7 +42,7 @@ TEST_FUNC constexpr bool test() int val = 42; using def_t = cuda::__argument::__deferred>; static_assert(cuda::__argument::__traits::lowest == 0); - static_assert(cuda::__argument::__traits::max == 100); + static_assert(cuda::__argument::__traits::highest == 100); // Also verify construction works auto def = cuda::__argument::__deferred{&val, cuda::__argument::__bounds<0, 100>()}; assert(cuda::__argument::__unwrap(def) == &val); @@ -54,7 +54,7 @@ TEST_FUNC constexpr bool test() auto def = cuda::__argument::__deferred{it, cuda::__argument::__bounds<0, 100>()}; assert(cuda::__argument::__unwrap(def)[0] == 42); static_assert(cuda::__argument::__traits::lowest == 0); - static_assert(cuda::__argument::__traits::max == 100); + static_assert(cuda::__argument::__traits::highest == 100); static_assert(cuda::__argument::__traits::is_single_value); } @@ -64,9 +64,9 @@ TEST_FUNC constexpr bool test() auto def = cuda::__argument::__deferred{ cuda::std::span{&val, 1}, cuda::__argument::__bounds(5, 100), cuda::__argument::__bounds<1, 256>()}; static_assert(cuda::__argument::__traits::lowest == 1); - static_assert(cuda::__argument::__traits::max == 256); + static_assert(cuda::__argument::__traits::highest == 256); assert(cuda::__argument::__lowest_(def) == 5); - assert(cuda::__argument::__max_(def) == 100); + assert(cuda::__argument::__highest_(def) == 100); } // Deferred sequence via fancy iterator @@ -76,7 +76,7 @@ TEST_FUNC constexpr bool test() assert(cuda::__argument::__unwrap(def)[0] == 10); assert(cuda::__argument::__unwrap(def)[2] == 12); static_assert(cuda::__argument::__traits::lowest == 0); - static_assert(cuda::__argument::__traits::max == 100); + static_assert(cuda::__argument::__traits::highest == 100); static_assert(!cuda::__argument::__traits::is_single_value); } @@ -87,7 +87,7 @@ TEST_FUNC constexpr bool test() cuda::std::span{arr, 4}, cuda::__argument::__bounds<1, 4096>(), cuda::__argument::__bounds(5, 100)}; static_assert(cuda::__argument::__traits::lowest == 1); assert(cuda::__argument::__lowest_(def) == 5); - assert(cuda::__argument::__max_(def) == 100); + assert(cuda::__argument::__highest_(def) == 100); } // Deferred sequence with both bounds, runtime bounds first @@ -96,9 +96,9 @@ TEST_FUNC constexpr bool test() auto def = cuda::__argument::__deferred_sequence{ cuda::std::span{arr, 4}, cuda::__argument::__bounds(5, 100), cuda::__argument::__bounds<1, 4096>()}; static_assert(cuda::__argument::__traits::lowest == 1); - static_assert(cuda::__argument::__traits::max == 4096); + static_assert(cuda::__argument::__traits::highest == 4096); assert(cuda::__argument::__lowest_(def) == 5); - assert(cuda::__argument::__max_(def) == 100); + assert(cuda::__argument::__highest_(def) == 100); } // Traits: deferred is single value diff --git a/libcudacxx/test/libcudacxx/cuda/argument/dynamic_argument.pass.cpp b/libcudacxx/test/libcudacxx/cuda/argument/dynamic_argument.pass.cpp index 8be2967f965..eac5e71aa82 100644 --- a/libcudacxx/test/libcudacxx/cuda/argument/dynamic_argument.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/argument/dynamic_argument.pass.cpp @@ -28,9 +28,9 @@ TEST_FUNC constexpr bool test() auto da = cuda::__argument::__immediate{5}; assert(cuda::__argument::__unwrap(da) == 5); static_assert(cuda::__argument::__traits::lowest == cuda::std::numeric_limits::lowest()); - static_assert(cuda::__argument::__traits::max == cuda::std::numeric_limits::max()); + static_assert(cuda::__argument::__traits::highest == (cuda::std::numeric_limits::max)()); assert(cuda::__argument::__lowest_(da) == 5); - assert(cuda::__argument::__max_(da) == 5); + assert(cuda::__argument::__highest_(da) == 5); } // Uniform scalar with static bounds @@ -38,9 +38,9 @@ TEST_FUNC constexpr bool test() auto da = cuda::__argument::__immediate{5, cuda::__argument::__bounds<1, 8>()}; assert(cuda::__argument::__unwrap(da) == 5); static_assert(cuda::__argument::__traits::lowest == 1); - static_assert(cuda::__argument::__traits::max == 8); + static_assert(cuda::__argument::__traits::highest == 8); assert(cuda::__argument::__lowest_(da) == 5); - assert(cuda::__argument::__max_(da) == 5); + assert(cuda::__argument::__highest_(da) == 5); } // Non-sequence values are accepted without scalar-only restrictions @@ -64,7 +64,7 @@ TEST_FUNC constexpr bool test() cuda::__argument::__immediate_sequence{cuda::std::span{arr, 4}, cuda::__argument::__bounds(1L, 100L)}; assert(cuda::__argument::__unwrap(da).size() == 4); assert(cuda::__argument::__lowest_(da) == 1); - assert(cuda::__argument::__max_(da) == 100); + assert(cuda::__argument::__highest_(da) == 100); } // Per-segment span with both bounds @@ -73,9 +73,9 @@ TEST_FUNC constexpr bool test() auto da = cuda::__argument::__immediate_sequence{ cuda::std::span{arr, 4}, cuda::__argument::__bounds<1, 256>(), cuda::__argument::__bounds(10, 200)}; static_assert(cuda::__argument::__traits::lowest == 1); - static_assert(cuda::__argument::__traits::max == 256); + static_assert(cuda::__argument::__traits::highest == 256); assert(cuda::__argument::__lowest_(da) == 10); - assert(cuda::__argument::__max_(da) == 200); + assert(cuda::__argument::__highest_(da) == 200); } // Per-segment span with both bounds, runtime bounds first @@ -84,9 +84,9 @@ TEST_FUNC constexpr bool test() auto da = cuda::__argument::__immediate_sequence{ cuda::std::span{arr, 4}, cuda::__argument::__bounds(10, 200), cuda::__argument::__bounds<1, 256>()}; static_assert(cuda::__argument::__traits::lowest == 1); - static_assert(cuda::__argument::__traits::max == 256); + static_assert(cuda::__argument::__traits::highest == 256); assert(cuda::__argument::__lowest_(da) == 10); - assert(cuda::__argument::__max_(da) == 200); + assert(cuda::__argument::__highest_(da) == 200); } // Per-segment via span @@ -106,7 +106,7 @@ TEST_FUNC constexpr bool test() assert(cuda::__argument::__unwrap(da).size() == 4); assert(cuda::__argument::__unwrap(da)[2] == 30); static_assert(cuda::__argument::__traits::lowest == 1); - static_assert(cuda::__argument::__traits::max == 100); + static_assert(cuda::__argument::__traits::highest == 100); } // Traits diff --git a/libcudacxx/test/libcudacxx/cuda/argument/static_argument.pass.cpp b/libcudacxx/test/libcudacxx/cuda/argument/static_argument.pass.cpp index 1495d59f7b4..7d40183070b 100644 --- a/libcudacxx/test/libcudacxx/cuda/argument/static_argument.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/argument/static_argument.pass.cpp @@ -66,15 +66,15 @@ TEST_FUNC void test() { constexpr auto sa = cuda::__argument::__constant<42>{}; static_assert(cuda::__argument::__lowest_(sa) == 42); - static_assert(cuda::__argument::__max_(sa) == 42); + static_assert(cuda::__argument::__highest_(sa) == 42); } #if TEST_HAS_CLASS_NTTP - // Bounds: array sequence computes min/max of elements + // Bounds: array sequence computes lowest/highest of elements { constexpr auto sa = cuda::__argument::__constant_sequence{128, 256, 512}>{}; static_assert(cuda::__argument::__lowest_(sa) == 128); - static_assert(cuda::__argument::__max_(sa) == 512); + static_assert(cuda::__argument::__highest_(sa) == 512); } #endif // TEST_HAS_CLASS_NTTP @@ -83,7 +83,7 @@ TEST_FUNC void test() { constexpr auto sa = cuda::__argument::__constant_sequence{}>{}; static_assert(cuda::__argument::__lowest_(sa) == cuda::std::numeric_limits::lowest()); - static_assert(cuda::__argument::__max_(sa) == cuda::std::numeric_limits::max()); + static_assert(cuda::__argument::__highest_(sa) == (cuda::std::numeric_limits::max)()); } #endif // TEST_HAS_CLASS_NTTP @@ -95,7 +95,7 @@ TEST_FUNC void test() static_assert(traits::is_single_value); static_assert(cuda::std::is_same_v); static_assert(traits::lowest == 42); - static_assert(traits::max == 42); + static_assert(traits::highest == 42); } #if TEST_HAS_CLASS_NTTP diff --git a/libcudacxx/test/libcudacxx/cuda/argument/static_bounds_conversion.fail.cpp b/libcudacxx/test/libcudacxx/cuda/argument/static_bounds_conversion.fail.cpp index 7767f5e7fe7..19e475ef453 100644 --- a/libcudacxx/test/libcudacxx/cuda/argument/static_bounds_conversion.fail.cpp +++ b/libcudacxx/test/libcudacxx/cuda/argument/static_bounds_conversion.fail.cpp @@ -12,7 +12,7 @@ using arg_t = cuda::__argument::__immediate>; -[[maybe_unused]] constexpr auto invalid_max = cuda::__argument::__traits::max; +[[maybe_unused]] constexpr auto invalid_highest = cuda::__argument::__traits::highest; int main(int, char**) { diff --git a/libcudacxx/test/libcudacxx/cuda/argument/usage_example.pass.cpp b/libcudacxx/test/libcudacxx/cuda/argument/usage_example.pass.cpp index 135d3355b11..e3a752b780a 100644 --- a/libcudacxx/test/libcudacxx/cuda/argument/usage_example.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/argument/usage_example.pass.cpp @@ -34,7 +34,7 @@ enum class algorithm_variant template TEST_FUNC constexpr algorithm_variant select_variant(_SegSizeArg) { - if constexpr (cuda::__argument::__traits<_SegSizeArg>::max <= shared_memory_capacity) + if constexpr (cuda::__argument::__traits<_SegSizeArg>::highest <= shared_memory_capacity) { return algorithm_variant::shared_memory; } @@ -48,8 +48,8 @@ TEST_FUNC constexpr algorithm_variant select_variant(_SegSizeArg) template TEST_FUNC constexpr int compute_buffer_size(_SegSizeArg __seg_size, int __num_segments) { - auto __max = cuda::std::min(default_max_segment_size, static_cast(cuda::__argument::__max_(__seg_size))); - return __max * __num_segments; + auto __highest = cuda::std::min(default_max_segment_size, static_cast(cuda::__argument::__highest_(__seg_size))); + return __highest * __num_segments; } // Process: use the actual unwrapped value. @@ -100,7 +100,7 @@ TEST_FUNC constexpr bool test() } #if TEST_HAS_CLASS_NTTP - // static_argument: array sequence, max fits in shared memory + // static_argument: array sequence, highest fits in shared memory { constexpr auto seg_sizes = cuda::__argument::__constant_sequence{}; static_assert(select_variant(seg_sizes) == algorithm_variant::shared_memory); @@ -108,7 +108,7 @@ TEST_FUNC constexpr bool test() assert(process_segments(seg_sizes) == 64 + 128 + 256); } - // static_argument: array sequence, max exceeds shared memory, buffer clamped + // static_argument: array sequence, highest exceeds shared memory, buffer clamped { constexpr auto seg_sizes = cuda::__argument::__constant_sequence{}; static_assert(select_variant(seg_sizes) == algorithm_variant::global_memory); @@ -156,7 +156,7 @@ TEST_FUNC constexpr bool test() int sizes[3] = {64, 128, 96}; auto seg_sizes = cuda::__argument::__immediate_sequence{ cuda::std::span{sizes, 3}, cuda::__argument::__bounds<1, 256>(), cuda::__argument::__bounds(1, 200)}; - static_assert(cuda::__argument::__traits::max <= shared_memory_capacity); + static_assert(cuda::__argument::__traits::highest <= shared_memory_capacity); assert(select_variant(seg_sizes) == algorithm_variant::shared_memory); assert(compute_buffer_size(seg_sizes, 3) == 200 * 3); assert(process_segments(seg_sizes) == 64 + 128 + 96); @@ -167,7 +167,7 @@ TEST_FUNC constexpr bool test() int val = 100; auto seg_size = cuda::__argument::__deferred{ cuda::std::span{&val, 1}, cuda::__argument::__bounds<1, 256>(), cuda::__argument::__bounds(1, 200)}; - static_assert(cuda::__argument::__traits::max <= shared_memory_capacity); + static_assert(cuda::__argument::__traits::highest <= shared_memory_capacity); assert(select_variant(seg_size) == algorithm_variant::shared_memory); assert(compute_buffer_size(seg_size, 4) == 200 * 4); }