Skip to content

Commit 6b7ae38

Browse files
pciolkoszpciolkosz
authored
[libcu++] Rename max to highest in arguments framework (#9246)
* Rename max to highest * Rebase
1 parent 6383b91 commit 6b7ae38

11 files changed

Lines changed: 118 additions & 116 deletions

File tree

‎cub/cub/agent/agent_batched_topk.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -95,7 +95,7 @@ struct agent_batched_topk_worker_per_segment
9595
multi_worker_per_segment_policy.threads_per_block * multi_worker_per_segment_policy.items_per_thread;
9696

9797
// Check if there could be large segments present
98-
static constexpr bool only_small_segments = ::cuda::__argument::__traits<SegmentSizeParameterT>::max <= tile_size;
98+
static constexpr bool only_small_segments = ::cuda::__argument::__traits<SegmentSizeParameterT>::highest <= tile_size;
9999

100100
// Check if we are dealing with keys-only or key-value pairs
101101
static constexpr bool is_keys_only = ::cuda::std::is_same_v<value_t, cub::NullType>;
@@ -224,7 +224,7 @@ struct agent_batched_topk_worker_per_segment
224224
const key_t padding_key =
225225
(direction == detail::topk::select::max)
226226
? ::cuda::std::numeric_limits<key_t>::lowest()
227-
: ::cuda::std::numeric_limits<key_t>::max();
227+
: (::cuda::std::numeric_limits<key_t>::max)();
228228

229229
// Dereference iterator-of-iterators to get the segment specific iterator
230230
auto block_keys_in = d_key_segments_it[segment_id];

‎cub/cub/device/dispatch/dispatch_batched_topk.cuh‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -126,7 +126,7 @@ template <typename KeyInputItItT,
126126
typename PolicySelector = policy_selector_from_types<it_value_t<it_value_t<KeyInputItItT>>,
127127
it_value_t<it_value_t<ValueInputItItT>>,
128128
::cuda::std::int64_t,
129-
::cuda::__argument::__traits<KParameterT>::max>>
129+
::cuda::__argument::__traits<KParameterT>::highest>>
130130
#if _CCCL_HAS_CONCEPTS()
131131
requires batched_topk_policy_selector<PolicySelector>
132132
#endif // _CCCL_HAS_CONCEPTS()
@@ -173,7 +173,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch(
173173
static constexpr bool any_small_segments =
174174
::cuda::__argument::__traits<SegmentSizeParameterT>::lowest <= worker_per_segment_tile_size;
175175
static constexpr bool only_small_segments =
176-
::cuda::__argument::__traits<SegmentSizeParameterT>::max <= worker_per_segment_tile_size;
176+
::cuda::__argument::__traits<SegmentSizeParameterT>::highest <= worker_per_segment_tile_size;
177177

178178
// Allocation layout:
179179
// only_small_segments: [0] dummy.
@@ -341,7 +341,7 @@ template <typename KeyInputItItT,
341341
policy_selector_from_types<it_value_t<it_value_t<KeyInputItItT>>,
342342
it_value_t<it_value_t<ValueInputItItT>>,
343343
::cuda::std::int64_t,
344-
::cuda::__argument::__traits<KParameterT>::max>;
344+
::cuda::__argument::__traits<KParameterT>::highest>;
345345
return detail::dispatch_with_env_and_tuning<default_policy_selector>(
346346
env, [&](auto policy_selector, void* d_temp_storage, size_t& temp_storage_bytes, cudaStream_t stream) {
347347
return dispatch(

‎cub/cub/device/dispatch/kernels/kernel_batched_topk.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -39,7 +39,7 @@ private:
3939
worker_policy worker_per_segment_policy;
4040
multi_worker_policy multi_worker_per_segment_policy;
4141
};
42-
static constexpr ::cuda::std::int64_t max_segment_size = ::cuda::__argument::__traits<SegmentSizeParameterT>::max;
42+
static constexpr ::cuda::std::int64_t max_segment_size = ::cuda::__argument::__traits<SegmentSizeParameterT>::highest;
4343
static constexpr batched_topk_policy active_policy = current_policy<PolicySelector>();
4444

4545
template <int Index>
@@ -151,7 +151,7 @@ __launch_bounds__(int(
151151
LargeSegmentTileOffsetT>::agent_t;
152152

153153
// Static Assertions (Constraints)
154-
static_assert(agent_t::tile_size >= ::cuda::__argument::__traits<SegmentSizeParameterT>::max,
154+
static_assert(agent_t::tile_size >= ::cuda::__argument::__traits<SegmentSizeParameterT>::highest,
155155
"Block size exceeds maximum segment size supported by SegmentSizeParameterT");
156156
static_assert(sizeof(typename agent_t::TempStorage) <= max_smem_per_block,
157157
"Static shared memory per block must not exceed 48KB limit.");

‎libcudacxx/include/cuda/__argument/argument.h‎

Lines changed: 67 additions & 66 deletions
Large diffs are not rendered by default.

‎libcudacxx/include/cuda/__argument/argument_bounds.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -77,7 +77,7 @@ template <class _Tp>
7777
struct __runtime_bounds
7878
{
7979
_Tp __lower_ = ::cuda::std::numeric_limits<_Tp>::lowest();
80-
_Tp __upper_ = ::cuda::std::numeric_limits<_Tp>::max();
80+
_Tp __upper_ = (::cuda::std::numeric_limits<_Tp>::max)();
8181

8282
constexpr __runtime_bounds() noexcept = default;
8383

‎libcudacxx/test/libcudacxx/cuda/argument/argument_traits.pass.cpp‎

Lines changed: 10 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -137,36 +137,37 @@ TEST_FUNC void test()
137137
cuda::std::array<int, 3>>);
138138
#endif // TEST_HAS_CLASS_NTTP
139139

140-
// --- argument_traits: lowest / max ---
140+
// --- argument_traits: lowest / highest ---
141141

142142
static_assert(cuda::__argument::__traits<int>::lowest == cuda::std::numeric_limits<int>::lowest());
143-
static_assert(cuda::__argument::__traits<int>::max == cuda::std::numeric_limits<int>::max());
143+
static_assert(cuda::__argument::__traits<int>::highest == (cuda::std::numeric_limits<int>::max)());
144144
static_assert(cuda::__argument::__traits<const int>::lowest == cuda::std::numeric_limits<int>::lowest());
145-
static_assert(cuda::__argument::__traits<int&>::max == cuda::std::numeric_limits<int>::max());
145+
static_assert(cuda::__argument::__traits<int&>::highest == (cuda::std::numeric_limits<int>::max)());
146146
static_assert(cuda::__argument::__traits<float>::lowest == cuda::std::numeric_limits<float>::lowest());
147-
static_assert(cuda::__argument::__traits<float>::max == cuda::std::numeric_limits<float>::max());
147+
static_assert(cuda::__argument::__traits<float>::highest == (cuda::std::numeric_limits<float>::max)());
148148
static_assert(
149149
cuda::__argument::__traits<const cuda::__argument::__immediate<int, cuda::__argument::__static_bounds<1, 8>>>::lowest
150150
== 1);
151151
static_assert(
152-
cuda::__argument::__traits<cuda::__argument::__immediate<int, cuda::__argument::__static_bounds<1, 8>>&>::max == 8);
152+
cuda::__argument::__traits<cuda::__argument::__immediate<int, cuda::__argument::__static_bounds<1, 8>>&>::highest
153+
== 8);
153154
static_assert(
154155
cuda::__argument::__traits<
155-
cuda::__argument::__immediate_sequence<cuda::std::span<int>, cuda::__argument::__static_bounds<1, 8>>>::max
156+
cuda::__argument::__immediate_sequence<cuda::std::span<int>, cuda::__argument::__static_bounds<1, 8>>>::highest
156157
== 8);
157158
#if TEST_HAS_CLASS_NTTP
158159
static_assert(
159160
cuda::__argument::__traits<cuda::__argument::__constant_sequence<cuda::std::array<int, 3>{3, 1, 2}>>::lowest == 1);
160161
static_assert(
161-
cuda::__argument::__traits<cuda::__argument::__constant_sequence<cuda::std::array<int, 3>{3, 1, 2}>>::max == 3);
162+
cuda::__argument::__traits<cuda::__argument::__constant_sequence<cuda::std::array<int, 3>{3, 1, 2}>>::highest == 3);
162163
#endif // TEST_HAS_CLASS_NTTP
163164

164165
// --- Free function bounds on plain values ---
165166

166167
static_assert(cuda::__argument::__lowest_(42) == cuda::std::numeric_limits<int>::lowest());
167-
static_assert(cuda::__argument::__max_(42) == cuda::std::numeric_limits<int>::max());
168+
static_assert(cuda::__argument::__highest_(42) == (cuda::std::numeric_limits<int>::max)());
168169
static_assert(cuda::__argument::__lowest_(1.0f) == cuda::std::numeric_limits<float>::lowest());
169-
static_assert(cuda::__argument::__max_(1.0f) == cuda::std::numeric_limits<float>::max());
170+
static_assert(cuda::__argument::__highest_(1.0f) == (cuda::std::numeric_limits<float>::max)());
170171

171172
// --- Scalar and sequence wrappers expose distinct single-value traits ---
172173

‎libcudacxx/test/libcudacxx/cuda/argument/deferred_argument.pass.cpp‎

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,7 @@ TEST_FUNC constexpr bool test()
2525
auto def = cuda::__argument::__deferred{cuda::std::span<int, 1>{&val, 1}};
2626
assert(cuda::__argument::__unwrap(def)[0] == 42);
2727
static_assert(cuda::__argument::__traits<decltype(def)>::lowest == cuda::std::numeric_limits<int>::lowest());
28-
static_assert(cuda::__argument::__traits<decltype(def)>::max == cuda::std::numeric_limits<int>::max());
28+
static_assert(cuda::__argument::__traits<decltype(def)>::highest == (cuda::std::numeric_limits<int>::max)());
2929
}
3030

3131
// Deferred single value with static bounds
@@ -34,15 +34,15 @@ TEST_FUNC constexpr bool test()
3434
auto def = cuda::__argument::__deferred{cuda::std::span<int, 1>{&val, 1}, cuda::__argument::__bounds<1, 1000>()};
3535
assert(cuda::__argument::__unwrap(def)[0] == 42);
3636
static_assert(cuda::__argument::__traits<decltype(def)>::lowest == 1);
37-
static_assert(cuda::__argument::__traits<decltype(def)>::max == 1000);
37+
static_assert(cuda::__argument::__traits<decltype(def)>::highest == 1000);
3838
}
3939

4040
// Deferred single value via pointer
4141
{
4242
int val = 42;
4343
using def_t = cuda::__argument::__deferred<int*, cuda::__argument::__static_bounds<0, 100>>;
4444
static_assert(cuda::__argument::__traits<def_t>::lowest == 0);
45-
static_assert(cuda::__argument::__traits<def_t>::max == 100);
45+
static_assert(cuda::__argument::__traits<def_t>::highest == 100);
4646
// Also verify construction works
4747
auto def = cuda::__argument::__deferred{&val, cuda::__argument::__bounds<0, 100>()};
4848
assert(cuda::__argument::__unwrap(def) == &val);
@@ -54,7 +54,7 @@ TEST_FUNC constexpr bool test()
5454
auto def = cuda::__argument::__deferred{it, cuda::__argument::__bounds<0, 100>()};
5555
assert(cuda::__argument::__unwrap(def)[0] == 42);
5656
static_assert(cuda::__argument::__traits<decltype(def)>::lowest == 0);
57-
static_assert(cuda::__argument::__traits<decltype(def)>::max == 100);
57+
static_assert(cuda::__argument::__traits<decltype(def)>::highest == 100);
5858
static_assert(cuda::__argument::__traits<decltype(def)>::is_single_value);
5959
}
6060

@@ -64,9 +64,9 @@ TEST_FUNC constexpr bool test()
6464
auto def = cuda::__argument::__deferred{
6565
cuda::std::span<int, 1>{&val, 1}, cuda::__argument::__bounds(5, 100), cuda::__argument::__bounds<1, 256>()};
6666
static_assert(cuda::__argument::__traits<decltype(def)>::lowest == 1);
67-
static_assert(cuda::__argument::__traits<decltype(def)>::max == 256);
67+
static_assert(cuda::__argument::__traits<decltype(def)>::highest == 256);
6868
assert(cuda::__argument::__lowest_(def) == 5);
69-
assert(cuda::__argument::__max_(def) == 100);
69+
assert(cuda::__argument::__highest_(def) == 100);
7070
}
7171

7272
// Deferred sequence via fancy iterator
@@ -76,7 +76,7 @@ TEST_FUNC constexpr bool test()
7676
assert(cuda::__argument::__unwrap(def)[0] == 10);
7777
assert(cuda::__argument::__unwrap(def)[2] == 12);
7878
static_assert(cuda::__argument::__traits<decltype(def)>::lowest == 0);
79-
static_assert(cuda::__argument::__traits<decltype(def)>::max == 100);
79+
static_assert(cuda::__argument::__traits<decltype(def)>::highest == 100);
8080
static_assert(!cuda::__argument::__traits<decltype(def)>::is_single_value);
8181
}
8282

@@ -87,7 +87,7 @@ TEST_FUNC constexpr bool test()
8787
cuda::std::span<int>{arr, 4}, cuda::__argument::__bounds<1, 4096>(), cuda::__argument::__bounds(5, 100)};
8888
static_assert(cuda::__argument::__traits<decltype(def)>::lowest == 1);
8989
assert(cuda::__argument::__lowest_(def) == 5);
90-
assert(cuda::__argument::__max_(def) == 100);
90+
assert(cuda::__argument::__highest_(def) == 100);
9191
}
9292

9393
// Deferred sequence with both bounds, runtime bounds first
@@ -96,9 +96,9 @@ TEST_FUNC constexpr bool test()
9696
auto def = cuda::__argument::__deferred_sequence{
9797
cuda::std::span<int>{arr, 4}, cuda::__argument::__bounds(5, 100), cuda::__argument::__bounds<1, 4096>()};
9898
static_assert(cuda::__argument::__traits<decltype(def)>::lowest == 1);
99-
static_assert(cuda::__argument::__traits<decltype(def)>::max == 4096);
99+
static_assert(cuda::__argument::__traits<decltype(def)>::highest == 4096);
100100
assert(cuda::__argument::__lowest_(def) == 5);
101-
assert(cuda::__argument::__max_(def) == 100);
101+
assert(cuda::__argument::__highest_(def) == 100);
102102
}
103103

104104
// Traits: deferred is single value

‎libcudacxx/test/libcudacxx/cuda/argument/dynamic_argument.pass.cpp‎

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -28,19 +28,19 @@ TEST_FUNC constexpr bool test()
2828
auto da = cuda::__argument::__immediate{5};
2929
assert(cuda::__argument::__unwrap(da) == 5);
3030
static_assert(cuda::__argument::__traits<decltype(da)>::lowest == cuda::std::numeric_limits<int>::lowest());
31-
static_assert(cuda::__argument::__traits<decltype(da)>::max == cuda::std::numeric_limits<int>::max());
31+
static_assert(cuda::__argument::__traits<decltype(da)>::highest == (cuda::std::numeric_limits<int>::max)());
3232
assert(cuda::__argument::__lowest_(da) == 5);
33-
assert(cuda::__argument::__max_(da) == 5);
33+
assert(cuda::__argument::__highest_(da) == 5);
3434
}
3535

3636
// Uniform scalar with static bounds
3737
{
3838
auto da = cuda::__argument::__immediate{5, cuda::__argument::__bounds<1, 8>()};
3939
assert(cuda::__argument::__unwrap(da) == 5);
4040
static_assert(cuda::__argument::__traits<decltype(da)>::lowest == 1);
41-
static_assert(cuda::__argument::__traits<decltype(da)>::max == 8);
41+
static_assert(cuda::__argument::__traits<decltype(da)>::highest == 8);
4242
assert(cuda::__argument::__lowest_(da) == 5);
43-
assert(cuda::__argument::__max_(da) == 5);
43+
assert(cuda::__argument::__highest_(da) == 5);
4444
}
4545

4646
// Non-sequence values are accepted without scalar-only restrictions
@@ -64,7 +64,7 @@ TEST_FUNC constexpr bool test()
6464
cuda::__argument::__immediate_sequence{cuda::std::span<int>{arr, 4}, cuda::__argument::__bounds(1L, 100L)};
6565
assert(cuda::__argument::__unwrap(da).size() == 4);
6666
assert(cuda::__argument::__lowest_(da) == 1);
67-
assert(cuda::__argument::__max_(da) == 100);
67+
assert(cuda::__argument::__highest_(da) == 100);
6868
}
6969

7070
// Per-segment span with both bounds
@@ -73,9 +73,9 @@ TEST_FUNC constexpr bool test()
7373
auto da = cuda::__argument::__immediate_sequence{
7474
cuda::std::span<int>{arr, 4}, cuda::__argument::__bounds<1, 256>(), cuda::__argument::__bounds(10, 200)};
7575
static_assert(cuda::__argument::__traits<decltype(da)>::lowest == 1);
76-
static_assert(cuda::__argument::__traits<decltype(da)>::max == 256);
76+
static_assert(cuda::__argument::__traits<decltype(da)>::highest == 256);
7777
assert(cuda::__argument::__lowest_(da) == 10);
78-
assert(cuda::__argument::__max_(da) == 200);
78+
assert(cuda::__argument::__highest_(da) == 200);
7979
}
8080

8181
// Per-segment span with both bounds, runtime bounds first
@@ -84,9 +84,9 @@ TEST_FUNC constexpr bool test()
8484
auto da = cuda::__argument::__immediate_sequence{
8585
cuda::std::span<int>{arr, 4}, cuda::__argument::__bounds(10, 200), cuda::__argument::__bounds<1, 256>()};
8686
static_assert(cuda::__argument::__traits<decltype(da)>::lowest == 1);
87-
static_assert(cuda::__argument::__traits<decltype(da)>::max == 256);
87+
static_assert(cuda::__argument::__traits<decltype(da)>::highest == 256);
8888
assert(cuda::__argument::__lowest_(da) == 10);
89-
assert(cuda::__argument::__max_(da) == 200);
89+
assert(cuda::__argument::__highest_(da) == 200);
9090
}
9191

9292
// Per-segment via span
@@ -106,7 +106,7 @@ TEST_FUNC constexpr bool test()
106106
assert(cuda::__argument::__unwrap(da).size() == 4);
107107
assert(cuda::__argument::__unwrap(da)[2] == 30);
108108
static_assert(cuda::__argument::__traits<decltype(da)>::lowest == 1);
109-
static_assert(cuda::__argument::__traits<decltype(da)>::max == 100);
109+
static_assert(cuda::__argument::__traits<decltype(da)>::highest == 100);
110110
}
111111

112112
// Traits

‎libcudacxx/test/libcudacxx/cuda/argument/static_argument.pass.cpp‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -66,15 +66,15 @@ TEST_FUNC void test()
6666
{
6767
constexpr auto sa = cuda::__argument::__constant<42>{};
6868
static_assert(cuda::__argument::__lowest_(sa) == 42);
69-
static_assert(cuda::__argument::__max_(sa) == 42);
69+
static_assert(cuda::__argument::__highest_(sa) == 42);
7070
}
7171

7272
#if TEST_HAS_CLASS_NTTP
73-
// Bounds: array sequence computes min/max of elements
73+
// Bounds: array sequence computes lowest/highest of elements
7474
{
7575
constexpr auto sa = cuda::__argument::__constant_sequence<cuda::std::array<int, 3>{128, 256, 512}>{};
7676
static_assert(cuda::__argument::__lowest_(sa) == 128);
77-
static_assert(cuda::__argument::__max_(sa) == 512);
77+
static_assert(cuda::__argument::__highest_(sa) == 512);
7878
}
7979
#endif // TEST_HAS_CLASS_NTTP
8080

@@ -83,7 +83,7 @@ TEST_FUNC void test()
8383
{
8484
constexpr auto sa = cuda::__argument::__constant_sequence<cuda::std::array<int, 0>{}>{};
8585
static_assert(cuda::__argument::__lowest_(sa) == cuda::std::numeric_limits<int>::lowest());
86-
static_assert(cuda::__argument::__max_(sa) == cuda::std::numeric_limits<int>::max());
86+
static_assert(cuda::__argument::__highest_(sa) == (cuda::std::numeric_limits<int>::max)());
8787
}
8888
#endif // TEST_HAS_CLASS_NTTP
8989

@@ -95,7 +95,7 @@ TEST_FUNC void test()
9595
static_assert(traits::is_single_value);
9696
static_assert(cuda::std::is_same_v<traits::value_type, int>);
9797
static_assert(traits::lowest == 42);
98-
static_assert(traits::max == 42);
98+
static_assert(traits::highest == 42);
9999
}
100100

101101
#if TEST_HAS_CLASS_NTTP

‎libcudacxx/test/libcudacxx/cuda/argument/static_bounds_conversion.fail.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,7 @@
1212

1313
using arg_t = cuda::__argument::__immediate<unsigned char, cuda::__argument::__static_bounds<0, 1000>>;
1414

15-
[[maybe_unused]] constexpr auto invalid_max = cuda::__argument::__traits<arg_t>::max;
15+
[[maybe_unused]] constexpr auto invalid_highest = cuda::__argument::__traits<arg_t>::highest;
1616

1717
int main(int, char**)
1818
{

0 commit comments

Comments
 (0)