Skip to content

Commit c66f76e

Browse files
griwesgriwes
committed
The great Thrust index type fix, part 7: partition point, reduce tests.
1 parent 9b164e2 commit c66f76e

File tree

4 files changed

+112
-57
lines changed

4 files changed

+112
-57
lines changed

‎testing/copy.cu‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -621,7 +621,7 @@ DECLARE_UNITTEST(TestCopyIfStencilDispatchImplicit);
621621

622622
struct only_set_when_expected_it
623623
{
624-
unsigned long long expected;
624+
long long expected;
625625
bool * flag;
626626

627627
__host__ __device__ only_set_when_expected_it operator++() const { return *this; }
@@ -653,14 +653,14 @@ struct iterator_traits<only_set_when_expected_it>
653653

654654
void TestCopyWithBigIndexesHelper(int magnitude)
655655
{
656-
thrust::counting_iterator<unsigned long long> begin(0);
657-
thrust::counting_iterator<unsigned long long> end = begin + (1ull << magnitude);
656+
thrust::counting_iterator<long long> begin(0);
657+
thrust::counting_iterator<long long> end = begin + (1ll << magnitude);
658658
ASSERT_EQUAL(thrust::distance(begin, end), 1ll << magnitude);
659659

660660
thrust::device_ptr<bool> has_executed = thrust::device_malloc<bool>(1);
661661
*has_executed = false;
662662

663-
only_set_when_expected_it out = { (1ull << magnitude) - 1, thrust::raw_pointer_cast(has_executed) };
663+
only_set_when_expected_it out = { (1ll << magnitude) - 1, thrust::raw_pointer_cast(has_executed) };
664664

665665
thrust::copy(thrust::device, begin, end, out);
666666

‎testing/partition_point.cu‎

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -95,3 +95,39 @@ void TestPartitionPointDispatchImplicit()
9595
}
9696
DECLARE_UNITTEST(TestPartitionPointDispatchImplicit);
9797

98+
struct test_less_than
99+
{
100+
long long expected;
101+
102+
__device__
103+
bool operator()(long long y)
104+
{
105+
return y < expected;
106+
}
107+
};
108+
109+
void TestPartitionPointWithBigIndexesHelper(int magnitude)
110+
{
111+
thrust::counting_iterator<long long> begin(0);
112+
thrust::counting_iterator<long long> end = begin + (1ll << magnitude);
113+
ASSERT_EQUAL(thrust::distance(begin, end), 1ll << magnitude);
114+
115+
test_less_than fn = { (1ll << magnitude) - 17 };
116+
117+
ASSERT_EQUAL(thrust::distance(
118+
begin,
119+
thrust::partition_point(
120+
thrust::device,
121+
begin, end,
122+
fn)),
123+
(1ll << magnitude) - 17);
124+
}
125+
126+
void TestPartitionPointWithBigIndexes()
127+
{
128+
TestPartitionPointWithBigIndexesHelper(30);
129+
TestPartitionPointWithBigIndexesHelper(31);
130+
TestPartitionPointWithBigIndexesHelper(32);
131+
TestPartitionPointWithBigIndexesHelper(33);
132+
}
133+
DECLARE_UNITTEST(TestPartitionPointWithBigIndexes);

‎testing/reduce.cu‎

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,7 @@
11
#include <unittest/unittest.h>
22
#include <thrust/reduce.h>
33
#include <thrust/iterator/counting_iterator.h>
4+
#include <thrust/iterator/constant_iterator.h>
45
#include <thrust/iterator/retag.h>
56
#include <limits>
67

@@ -210,3 +211,22 @@ template<typename T>
210211
}
211212
DECLARE_GENERIC_UNITTEST(TestReduceCountingIterator);
212213

214+
void TestReduceWithBigIndexesHelper(int magnitude)
215+
{
216+
thrust::constant_iterator<long long> begin(1);
217+
thrust::constant_iterator<long long> end = begin + (1ll << magnitude);
218+
ASSERT_EQUAL(thrust::distance(begin, end), 1ll << magnitude);
219+
220+
long long result = thrust::reduce(thrust::device, begin, end);
221+
222+
ASSERT_EQUAL(result, 1ll << magnitude);
223+
}
224+
225+
void TestReduceWithBigIndexes()
226+
{
227+
TestReduceWithBigIndexesHelper(30);
228+
TestReduceWithBigIndexesHelper(31);
229+
TestReduceWithBigIndexesHelper(32);
230+
TestReduceWithBigIndexesHelper(33);
231+
}
232+
DECLARE_UNITTEST(TestReduceWithBigIndexes);

‎testing/remove.cu‎

Lines changed: 52 additions & 53 deletions
Original file line numberDiff line numberDiff line change
@@ -30,14 +30,14 @@ void TestRemoveSimple(void)
3030
typedef typename Vector::value_type T;
3131

3232
Vector data(5);
33-
data[0] = 1;
34-
data[1] = 2;
33+
data[0] = 1;
34+
data[1] = 2;
3535
data[2] = 1;
36-
data[3] = 3;
37-
data[4] = 2;
36+
data[3] = 3;
37+
data[4] = 2;
3838

39-
typename Vector::iterator end = thrust::remove(data.begin(),
40-
data.end(),
39+
typename Vector::iterator end = thrust::remove(data.begin(),
40+
data.end(),
4141
(T) 2);
4242

4343
ASSERT_EQUAL(end - data.begin(), 3);
@@ -102,17 +102,17 @@ void TestRemoveCopySimple(void)
102102
typedef typename Vector::value_type T;
103103

104104
Vector data(5);
105-
data[0] = 1;
106-
data[1] = 2;
105+
data[0] = 1;
106+
data[1] = 2;
107107
data[2] = 1;
108-
data[3] = 3;
109-
data[4] = 2;
108+
data[3] = 3;
109+
data[4] = 2;
110110

111111
Vector result(5);
112112

113-
typename Vector::iterator end = thrust::remove_copy(data.begin(),
114-
data.end(),
115-
result.begin(),
113+
typename Vector::iterator end = thrust::remove_copy(data.begin(),
114+
data.end(),
115+
result.begin(),
116116
(T) 2);
117117

118118
ASSERT_EQUAL(end - result.begin(), 3);
@@ -186,14 +186,14 @@ void TestRemoveIfSimple(void)
186186
typedef typename Vector::value_type T;
187187

188188
Vector data(5);
189-
data[0] = 1;
190-
data[1] = 2;
189+
data[0] = 1;
190+
data[1] = 2;
191191
data[2] = 1;
192-
data[3] = 3;
193-
data[4] = 2;
192+
data[3] = 3;
193+
data[4] = 2;
194194

195-
typename Vector::iterator end = thrust::remove_if(data.begin(),
196-
data.end(),
195+
typename Vector::iterator end = thrust::remove_if(data.begin(),
196+
data.end(),
197197
is_even<T>());
198198

199199
ASSERT_EQUAL(end - data.begin(), 3);
@@ -258,11 +258,11 @@ void TestRemoveIfStencilSimple(void)
258258
typedef typename Vector::value_type T;
259259

260260
Vector data(5);
261-
data[0] = 1;
262-
data[1] = 2;
261+
data[0] = 1;
262+
data[1] = 2;
263263
data[2] = 1;
264-
data[3] = 3;
265-
data[4] = 2;
264+
data[3] = 3;
265+
data[4] = 2;
266266

267267
Vector stencil(5);
268268
stencil[0] = 0;
@@ -271,7 +271,7 @@ void TestRemoveIfStencilSimple(void)
271271
stencil[3] = 0;
272272
stencil[4] = 1;
273273

274-
typename Vector::iterator end = thrust::remove_if(data.begin(),
274+
typename Vector::iterator end = thrust::remove_if(data.begin(),
275275
data.end(),
276276
stencil.begin(),
277277
thrust::identity<T>());
@@ -347,17 +347,17 @@ void TestRemoveCopyIfSimple(void)
347347
typedef typename Vector::value_type T;
348348

349349
Vector data(5);
350-
data[0] = 1;
351-
data[1] = 2;
350+
data[0] = 1;
351+
data[1] = 2;
352352
data[2] = 1;
353-
data[3] = 3;
354-
data[4] = 2;
353+
data[3] = 3;
354+
data[4] = 2;
355355

356356
Vector result(5);
357357

358-
typename Vector::iterator end = thrust::remove_copy_if(data.begin(),
359-
data.end(),
360-
result.begin(),
358+
typename Vector::iterator end = thrust::remove_copy_if(data.begin(),
359+
data.end(),
360+
result.begin(),
361361
is_even<T>());
362362

363363
ASSERT_EQUAL(end - result.begin(), 3);
@@ -431,11 +431,11 @@ void TestRemoveCopyIfStencilSimple(void)
431431
typedef typename Vector::value_type T;
432432

433433
Vector data(5);
434-
data[0] = 1;
435-
data[1] = 2;
434+
data[0] = 1;
435+
data[1] = 2;
436436
data[2] = 1;
437-
data[3] = 3;
438-
data[4] = 2;
437+
data[3] = 3;
438+
data[4] = 2;
439439

440440
Vector stencil(5);
441441
stencil[0] = 0;
@@ -446,10 +446,10 @@ void TestRemoveCopyIfStencilSimple(void)
446446

447447
Vector result(5);
448448

449-
typename Vector::iterator end = thrust::remove_copy_if(data.begin(),
450-
data.end(),
449+
typename Vector::iterator end = thrust::remove_copy_if(data.begin(),
450+
data.end(),
451451
stencil.begin(),
452-
result.begin(),
452+
result.begin(),
453453
thrust::identity<T>());
454454

455455
ASSERT_EQUAL(end - result.begin(), 3);
@@ -531,7 +531,7 @@ void TestRemove(const size_t n)
531531

532532
size_t h_size = thrust::remove(h_data.begin(), h_data.end(), T(0)) - h_data.begin();
533533
size_t d_size = thrust::remove(d_data.begin(), d_data.end(), T(0)) - d_data.begin();
534-
534+
535535
ASSERT_EQUAL(h_size, d_size);
536536

537537
h_data.resize(h_size);
@@ -550,7 +550,7 @@ void TestRemoveIf(const size_t n)
550550

551551
size_t h_size = thrust::remove_if(h_data.begin(), h_data.end(), is_true<T>()) - h_data.begin();
552552
size_t d_size = thrust::remove_if(d_data.begin(), d_data.end(), is_true<T>()) - d_data.begin();
553-
553+
554554
ASSERT_EQUAL(h_size, d_size);
555555

556556
h_data.resize(h_size);
@@ -569,10 +569,10 @@ void TestRemoveIfStencil(const size_t n)
569569

570570
thrust::host_vector<bool> h_stencil = unittest::random_integers<bool>(n);
571571
thrust::device_vector<bool> d_stencil = h_stencil;
572-
572+
573573
size_t h_size = thrust::remove_if(h_data.begin(), h_data.end(), h_stencil.begin(), is_true<T>()) - h_data.begin();
574574
size_t d_size = thrust::remove_if(d_data.begin(), d_data.end(), d_stencil.begin(), is_true<T>()) - d_data.begin();
575-
575+
576576
ASSERT_EQUAL(h_size, d_size);
577577

578578
h_data.resize(h_size);
@@ -588,13 +588,13 @@ void TestRemoveCopy(const size_t n)
588588
{
589589
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
590590
thrust::device_vector<T> d_data = h_data;
591-
591+
592592
thrust::host_vector<T> h_result(n);
593593
thrust::device_vector<T> d_result(n);
594594

595595
size_t h_size = thrust::remove_copy(h_data.begin(), h_data.end(), h_result.begin(), T(0)) - h_result.begin();
596596
size_t d_size = thrust::remove_copy(d_data.begin(), d_data.end(), d_result.begin(), T(0)) - d_result.begin();
597-
597+
598598
ASSERT_EQUAL(h_size, d_size);
599599

600600
h_result.resize(h_size);
@@ -621,7 +621,7 @@ void TestRemoveCopyToDiscardIterator(const size_t n)
621621
thrust::remove_copy(d_data.begin(), d_data.end(), thrust::make_discard_iterator(), T(0));
622622

623623
thrust::discard_iterator<> reference(num_nonzeros);
624-
624+
625625
ASSERT_EQUAL_QUIET(reference, h_result);
626626
ASSERT_EQUAL_QUIET(reference, d_result);
627627
}
@@ -659,7 +659,7 @@ void TestRemoveCopyToDiscardIteratorZipped(const size_t n)
659659
thrust::make_tuple(T(0),T(0)));
660660

661661
thrust::discard_iterator<> reference(num_nonzeros);
662-
662+
663663
ASSERT_EQUAL(h_output, d_output);
664664
ASSERT_EQUAL_QUIET(reference, thrust::get<1>(h_result.get_iterator_tuple()));
665665
ASSERT_EQUAL_QUIET(reference, thrust::get<1>(d_result.get_iterator_tuple()));
@@ -675,10 +675,10 @@ void TestRemoveCopyIf(const size_t n)
675675

676676
thrust::host_vector<T> h_result(n);
677677
thrust::device_vector<T> d_result(n);
678-
678+
679679
size_t h_size = thrust::remove_copy_if(h_data.begin(), h_data.end(), h_result.begin(), is_true<T>()) - h_result.begin();
680680
size_t d_size = thrust::remove_copy_if(d_data.begin(), d_data.end(), d_result.begin(), is_true<T>()) - d_result.begin();
681-
681+
682682
ASSERT_EQUAL(h_size, d_size);
683683

684684
h_result.resize(h_size);
@@ -716,16 +716,16 @@ void TestRemoveCopyIfStencil(const size_t n)
716716
{
717717
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
718718
thrust::device_vector<T> d_data = h_data;
719-
719+
720720
thrust::host_vector<bool> h_stencil = unittest::random_integers<bool>(n);
721721
thrust::device_vector<bool> d_stencil = h_stencil;
722-
722+
723723
thrust::host_vector<T> h_result(n);
724724
thrust::device_vector<T> d_result(n);
725725

726726
size_t h_size = thrust::remove_copy_if(h_data.begin(), h_data.end(), h_stencil.begin(), h_result.begin(), is_true<T>()) - h_result.begin();
727727
size_t d_size = thrust::remove_copy_if(d_data.begin(), d_data.end(), d_stencil.begin(), d_result.begin(), is_true<T>()) - d_result.begin();
728-
728+
729729
ASSERT_EQUAL(h_size, d_size);
730730

731731
h_result.resize(h_size);
@@ -741,7 +741,7 @@ void TestRemoveCopyIfStencilToDiscardIterator(const size_t n)
741741
{
742742
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
743743
thrust::device_vector<T> d_data = h_data;
744-
744+
745745
thrust::host_vector<bool> h_stencil = unittest::random_integers<bool>(n);
746746
thrust::device_vector<bool> d_stencil = h_stencil;
747747

@@ -759,4 +759,3 @@ void TestRemoveCopyIfStencilToDiscardIterator(const size_t n)
759759
ASSERT_EQUAL_QUIET(reference, d_result);
760760
}
761761
DECLARE_VARIABLE_UNITTEST(TestRemoveCopyIfStencilToDiscardIterator);
762-

0 commit comments

Comments
 (0)