Skip to content

Commit c339a52

Browse files
bernhardmgruberbernhardmgruber
authored
Use cuda::std::min/max in Thrust (#3364)
1 parent 6a0f48b commit c339a52

File tree

24 files changed

+96
-308
lines changed

24 files changed

+96
-308
lines changed

‎cub/test/catch2_large_array_sort_helper.cuh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -140,7 +140,7 @@ public:
140140
_CCCL_HOST_DEVICE KeyType operator()(std::size_t idx) const
141141
{
142142
// The final summary may be padded, so truncate the summary_idx at the last valid idx:
143-
const std::size_t summary_idx = thrust::min(m_num_summaries - 1, idx / m_unpadded_run_size);
143+
const std::size_t summary_idx = cuda::std::min(m_num_summaries - 1, idx / m_unpadded_run_size);
144144
const KeyType key = m_is_descending ? static_cast<KeyType>((m_num_summaries - 1 - summary_idx) * m_key_conversion)
145145
: static_cast<KeyType>(summary_idx * m_key_conversion);
146146

‎thrust/examples/set_operations.cu‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,5 @@
11
#include <thrust/device_vector.h>
2+
#include <thrust/extrema.h>
23
#include <thrust/iterator/discard_iterator.h>
34
#include <thrust/merge.h>
45
#include <thrust/set_operations.h>

‎thrust/testing/async/exclusive_scan/large_indices.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -196,17 +196,17 @@ namespace
196196
{
197197

198198
//------------------------------------------------------------------------------
199-
// Generate the output sequence using counting iterators and thrust::max<> for
199+
// Generate the output sequence using counting iterators and ::cuda::maximum<> for
200200
// custom operator overloads.
201201
struct custom_bin_op_overloads
202202
{
203203
using postfix_args_type = std::tuple< // List any extra arg overloads:
204-
std::tuple<uint64_t, thrust::maximum<>> // - initial_value, binop
204+
std::tuple<uint64_t, ::cuda::maximum<>> // - initial_value, binop
205205
>;
206206

207207
static postfix_args_type generate_postfix_args()
208208
{
209-
return postfix_args_type{std::make_tuple(0, thrust::maximum<>{})};
209+
return postfix_args_type{std::make_tuple(0, ::cuda::maximum<>{})};
210210
}
211211
};
212212

‎thrust/testing/async/inclusive_scan/large_indices.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -191,17 +191,17 @@ namespace
191191
{
192192

193193
//------------------------------------------------------------------------------
194-
// Generate the output sequence using counting iterators and thrust::max<> for
194+
// Generate the output sequence using counting iterators and ::cuda::maximum<> for
195195
// custom operator overloads.
196196
struct custom_bin_op_overloads
197197
{
198198
using postfix_args_type = std::tuple< // List any extra arg overloads:
199-
std::tuple<thrust::maximum<>> // - custom binary op
199+
std::tuple<::cuda::maximum<>> // - custom binary op
200200
>;
201201

202202
static postfix_args_type generate_postfix_args()
203203
{
204-
return postfix_args_type{std::make_tuple(thrust::maximum<>{})};
204+
return postfix_args_type{std::make_tuple(::cuda::maximum<>{})};
205205
}
206206
};
207207

‎thrust/testing/cuda/is_partitioned.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,7 @@ void TestIsPartitionedDevice(ExecutionPolicy exec)
2626
{
2727
size_t n = 1000;
2828

29-
n = thrust::max<size_t>(n, 2);
29+
n = ::cuda::std::max<size_t>(n, 2);
3030

3131
thrust::device_vector<int> v = unittest::random_integers<int>(n);
3232

‎thrust/testing/min_and_max.cu‎

Lines changed: 24 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -9,32 +9,28 @@ struct TestMin
99
{
1010
// 2 < 3
1111
T two(2), three(3);
12-
ASSERT_EQUAL(two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two, three));
13-
ASSERT_EQUAL(two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two, three, thrust::less<T>()));
12+
ASSERT_EQUAL(two, ::cuda::std::min(two, three));
13+
ASSERT_EQUAL(two, ::cuda::std::min(two, three, thrust::less<T>()));
1414

15-
ASSERT_EQUAL(two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(three, two));
16-
ASSERT_EQUAL(two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::less<T>()));
15+
ASSERT_EQUAL(two, ::cuda::std::min(three, two));
16+
ASSERT_EQUAL(two, ::cuda::std::min(three, two, thrust::less<T>()));
1717

18-
ASSERT_EQUAL(three, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two, three, thrust::greater<T>()));
19-
ASSERT_EQUAL(three, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::greater<T>()));
18+
ASSERT_EQUAL(three, ::cuda::std::min(two, three, thrust::greater<T>()));
19+
ASSERT_EQUAL(three, ::cuda::std::min(three, two, thrust::greater<T>()));
2020

2121
using KV = key_value<T, T>;
2222
KV two_and_two(two, two);
2323
KV two_and_three(two, three);
2424

2525
// the first element breaks ties
26-
ASSERT_EQUAL_QUIET(two_and_two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three));
27-
ASSERT_EQUAL_QUIET(two_and_three, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two));
26+
ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::min(two_and_two, two_and_three));
27+
ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::min(two_and_three, two_and_two));
2828

29-
ASSERT_EQUAL_QUIET(two_and_two,
30-
thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three, thrust::less<KV>()));
31-
ASSERT_EQUAL_QUIET(two_and_three,
32-
thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::less<KV>()));
29+
ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::min(two_and_two, two_and_three, thrust::less<KV>()));
30+
ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::min(two_and_three, two_and_two, thrust::less<KV>()));
3331

34-
ASSERT_EQUAL_QUIET(
35-
two_and_two, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three, thrust::greater<KV>()));
36-
ASSERT_EQUAL_QUIET(
37-
two_and_three, thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::greater<KV>()));
32+
ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::min(two_and_two, two_and_three, thrust::greater<KV>()));
33+
ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::min(two_and_three, two_and_two, thrust::greater<KV>()));
3834
}
3935
};
4036
SimpleUnitTest<TestMin, NumericTypes> TestMinInstance;
@@ -46,32 +42,28 @@ struct TestMax
4642
{
4743
// 2 < 3
4844
T two(2), three(3);
49-
ASSERT_EQUAL(three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two, three));
50-
ASSERT_EQUAL(three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two, three, thrust::less<T>()));
45+
ASSERT_EQUAL(three, ::cuda::std::max(two, three));
46+
ASSERT_EQUAL(three, ::cuda::std::max(two, three, thrust::less<T>()));
5147

52-
ASSERT_EQUAL(three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(three, two));
53-
ASSERT_EQUAL(three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::less<T>()));
48+
ASSERT_EQUAL(three, ::cuda::std::max(three, two));
49+
ASSERT_EQUAL(three, ::cuda::std::max(three, two, thrust::less<T>()));
5450

55-
ASSERT_EQUAL(two, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two, three, thrust::greater<T>()));
56-
ASSERT_EQUAL(two, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(three, two, thrust::greater<T>()));
51+
ASSERT_EQUAL(two, ::cuda::std::max(two, three, thrust::greater<T>()));
52+
ASSERT_EQUAL(two, ::cuda::std::max(three, two, thrust::greater<T>()));
5753

5854
using KV = key_value<T, T>;
5955
KV two_and_two(two, two);
6056
KV two_and_three(two, three);
6157

6258
// the first element breaks ties
63-
ASSERT_EQUAL_QUIET(two_and_two, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three));
64-
ASSERT_EQUAL_QUIET(two_and_three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two));
59+
ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::max(two_and_two, two_and_three));
60+
ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::max(two_and_three, two_and_two));
6561

66-
ASSERT_EQUAL_QUIET(two_and_two,
67-
thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three, thrust::less<KV>()));
68-
ASSERT_EQUAL_QUIET(two_and_three,
69-
thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::less<KV>()));
62+
ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::max(two_and_two, two_and_three, thrust::less<KV>()));
63+
ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::max(two_and_three, two_and_two, thrust::less<KV>()));
7064

71-
ASSERT_EQUAL_QUIET(
72-
two_and_two, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_two, two_and_three, thrust::greater<KV>()));
73-
ASSERT_EQUAL_QUIET(
74-
two_and_three, thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(two_and_three, two_and_two, thrust::greater<KV>()));
65+
ASSERT_EQUAL_QUIET(two_and_two, ::cuda::std::max(two_and_two, two_and_three, thrust::greater<KV>()));
66+
ASSERT_EQUAL_QUIET(two_and_three, ::cuda::std::max(two_and_three, two_and_two, thrust::greater<KV>()));
7567
}
7668
};
7769
SimpleUnitTest<TestMax, NumericTypes> TestMaxInstance;

‎thrust/testing/scan.cu‎

Lines changed: 11 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -8,19 +8,11 @@
88
#include <thrust/iterator/retag.h>
99
#include <thrust/scan.h>
1010

11+
#include <cuda/functional>
1112
#include <cuda/std/array>
1213

1314
#include <unittest/unittest.h>
1415

15-
template <typename T>
16-
struct max_functor
17-
{
18-
_CCCL_HOST_DEVICE T operator()(T rhs, T lhs) const
19-
{
20-
return thrust::max(rhs, lhs);
21-
}
22-
};
23-
2416
template <class Vector>
2517
void TestScanSimple()
2618
{
@@ -289,12 +281,12 @@ struct TestScanWithOperator
289281
thrust::host_vector<T> h_output(n);
290282
thrust::device_vector<T> d_output(n);
291283

292-
thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), max_functor<T>());
293-
thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), max_functor<T>());
284+
thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), cuda::maximum<T>{});
285+
thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), cuda::maximum<T>{});
294286
ASSERT_EQUAL(d_output, h_output);
295287

296-
thrust::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), T(13), max_functor<T>());
297-
thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), T(13), max_functor<T>());
288+
thrust::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), T(13), cuda::maximum<T>{});
289+
thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), T(13), cuda::maximum<T>{});
298290
ASSERT_EQUAL(d_output, h_output);
299291
}
300292
};
@@ -311,19 +303,19 @@ struct TestScanWithOperatorToDiscardIterator
311303
thrust::discard_iterator<> reference(n);
312304

313305
thrust::discard_iterator<> h_result =
314-
thrust::inclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), max_functor<T>());
306+
thrust::inclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), cuda::maximum<T>{});
315307

316308
thrust::discard_iterator<> d_result =
317-
thrust::inclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), max_functor<T>());
309+
thrust::inclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), cuda::maximum<T>{});
318310

319311
ASSERT_EQUAL_QUIET(reference, h_result);
320312
ASSERT_EQUAL_QUIET(reference, d_result);
321313

322-
h_result =
323-
thrust::exclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), T(13), max_functor<T>());
314+
h_result = thrust::exclusive_scan(
315+
h_input.begin(), h_input.end(), thrust::make_discard_iterator(), T(13), cuda::maximum<T>{});
324316

325-
d_result =
326-
thrust::exclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), T(13), max_functor<T>());
317+
d_result = thrust::exclusive_scan(
318+
d_input.begin(), d_input.end(), thrust::make_discard_iterator(), T(13), cuda::maximum<T>{});
327319

328320
ASSERT_EQUAL_QUIET(reference, h_result);
329321
ASSERT_EQUAL_QUIET(reference, d_result);

‎thrust/thrust/detail/minmax.h‎

Lines changed: 0 additions & 55 deletions
This file was deleted.

‎thrust/thrust/detail/vector_base.inl‎

Lines changed: 16 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,6 @@
2727
#endif // no system header
2828
#include <thrust/advance.h>
2929
#include <thrust/detail/copy.h>
30-
#include <thrust/detail/minmax.h>
3130
#include <thrust/detail/overlapped_copy.h>
3231
#include <thrust/detail/temporary_array.h>
3332
#include <thrust/detail/type_traits.h>
@@ -36,6 +35,9 @@
3635
#include <thrust/equal.h>
3736
#include <thrust/iterator/iterator_traits.h>
3837

38+
#include <cuda/std/__algorithm/max.h>
39+
#include <cuda/std/__algorithm/min.h>
40+
3941
#include <stdexcept>
4042

4143
THRUST_NAMESPACE_BEGIN
@@ -348,7 +350,7 @@ void vector_base<T, Alloc>::reserve(size_type n)
348350
size_type new_capacity = n;
349351

350352
// do not exceed maximum storage
351-
new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION<size_type>(new_capacity, max_size());
353+
new_capacity = ::cuda::std::min<size_type>(new_capacity, max_size());
352354

353355
// create new storage
354356
storage_type new_storage(copy_allocator_t(), m_storage, new_capacity);
@@ -726,13 +728,14 @@ void vector_base<T, Alloc>::copy_insert(iterator position, ForwardIterator first
726728
const size_type old_size = size();
727729

728730
// compute the new capacity after the allocation
729-
size_type new_capacity = old_size + thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, num_new_elements);
731+
size_type new_capacity =
732+
old_size + ::cuda::std::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, num_new_elements);
730733

731734
// allocate exponentially larger new storage
732-
new_capacity = thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION<size_type>(new_capacity, 2 * capacity());
735+
new_capacity = ::cuda::std::max<size_type>(new_capacity, 2 * capacity());
733736

734737
// do not exceed maximum storage
735-
new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION<size_type>(new_capacity, max_size());
738+
new_capacity = ::cuda::std::min<size_type>(new_capacity, max_size());
736739

737740
if (new_capacity > max_size())
738741
{
@@ -797,13 +800,13 @@ void vector_base<T, Alloc>::append(size_type n)
797800
const size_type old_size = size();
798801

799802
// compute the new capacity after the allocation
800-
size_type new_capacity = old_size + thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, n);
803+
size_type new_capacity = old_size + ::cuda::std::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, n);
801804

802805
// allocate exponentially larger new storage
803-
new_capacity = thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION<size_type>(new_capacity, 2 * capacity());
806+
new_capacity = ::cuda::std::max<size_type>(new_capacity, 2 * capacity());
804807

805808
// do not exceed maximum storage
806-
new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION<size_type>(new_capacity, max_size());
809+
new_capacity = ::cuda::std::min<size_type>(new_capacity, max_size());
807810

808811
// create new storage
809812
storage_type new_storage(copy_allocator_t(), m_storage, new_capacity);
@@ -892,13 +895,13 @@ void vector_base<T, Alloc>::fill_insert(iterator position, size_type n, const T&
892895
const size_type old_size = size();
893896

894897
// compute the new capacity after the allocation
895-
size_type new_capacity = old_size + thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, n);
898+
size_type new_capacity = old_size + ::cuda::std::max THRUST_PREVENT_MACRO_SUBSTITUTION(old_size, n);
896899

897900
// allocate exponentially larger new storage
898-
new_capacity = thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION<size_type>(new_capacity, 2 * capacity());
901+
new_capacity = ::cuda::std::max<size_type>(new_capacity, 2 * capacity());
899902

900903
// do not exceed maximum storage
901-
new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION<size_type>(new_capacity, max_size());
904+
new_capacity = ::cuda::std::min<size_type>(new_capacity, max_size());
902905

903906
if (new_capacity > max_size())
904907
{
@@ -1072,10 +1075,10 @@ void vector_base<T, Alloc>::allocate_and_copy(
10721075
} // end if
10731076

10741077
// allocate exponentially larger new storage
1075-
size_type allocated_size = thrust::max<size_type>(requested_size, 2 * capacity());
1078+
size_type allocated_size = ::cuda::std::max<size_type>(requested_size, 2 * capacity());
10761079

10771080
// do not exceed maximum storage
1078-
allocated_size = thrust::min<size_type>(allocated_size, max_size());
1081+
allocated_size = ::cuda::std::min<size_type>(allocated_size, max_size());
10791082

10801083
if (requested_size > allocated_size)
10811084
{

0 commit comments

Comments
 (0)