The great Thrust index type fix, part 7: partition point, reduce tests.
MODIFY
testing/copy.cu
testing/copy.cu
@@ -653,7 +653,7 @@ DECLARE_UNITTEST(TestCopyIfStencilDispatchImplicit);
struct only_set_when_expected_it
{
- unsigned long long expected;
+ long long expected;
bool * flag;
__host__ __device__ only_set_when_expected_it operator++() const { return *this; }
@@ -685,14 +685,14 @@ struct iterator_traits<only_set_when_expected_it>
void TestCopyWithBigIndexesHelper(int magnitude)
{
- thrust::counting_iterator<unsigned long long> begin(0);
- thrust::counting_iterator<unsigned long long> end = begin + (1ull << magnitude);
+ thrust::counting_iterator<long long> begin(0);
+ thrust::counting_iterator<long long> end = begin + (1ll << magnitude);
ASSERT_EQUAL(thrust::distance(begin, end), 1ll << magnitude);
thrust::device_ptr<bool> has_executed = thrust::device_malloc<bool>(1);
*has_executed = false;
- only_set_when_expected_it out = { (1ull << magnitude) - 1, thrust::raw_pointer_cast(has_executed) };
+ only_set_when_expected_it out = { (1ll << magnitude) - 1, thrust::raw_pointer_cast(has_executed) };
thrust::copy(thrust::device, begin, end, out);

MODIFY
testing/partition_point.cu
testing/partition_point.cu
@@ -113,3 +113,40 @@ void TestPartitionPointDispatchImplicit()
ASSERT_EQUAL(13, vec.front());
}
DECLARE_UNITTEST(TestPartitionPointDispatchImplicit);
+
+struct test_less_than
+{
+ long long expected;
+
+ __device__
+ bool operator()(long long y)
+ {
+ return y < expected;
+ }
+};
+
+void TestPartitionPointWithBigIndexesHelper(int magnitude)
+{
+ thrust::counting_iterator<long long> begin(0);
+ thrust::counting_iterator<long long> end = begin + (1ll << magnitude);
+ ASSERT_EQUAL(thrust::distance(begin, end), 1ll << magnitude);
+
+ test_less_than fn = { (1ll << magnitude) - 17 };
+
+ ASSERT_EQUAL(thrust::distance(
+ begin,
+ thrust::partition_point(
+ thrust::device,
+ begin, end,
+ fn)),
+ (1ll << magnitude) - 17);
+}
+
+void TestPartitionPointWithBigIndexes()
+{
+ TestPartitionPointWithBigIndexesHelper(30);
+ TestPartitionPointWithBigIndexesHelper(31);
+ TestPartitionPointWithBigIndexesHelper(32);
+ TestPartitionPointWithBigIndexesHelper(33);
+}
+DECLARE_UNITTEST(TestPartitionPointWithBigIndexes);

MODIFY
testing/reduce.cu
testing/reduce.cu
#include <unittest/unittest.h>
#include <thrust/reduce.h>
#include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/constant_iterator.h>
#include <thrust/iterator/retag.h>
#include <limits>
@@ -210,3 +211,22 @@ template<typename T>
}
DECLARE_GENERIC_UNITTEST(TestReduceCountingIterator);
+void TestReduceWithBigIndexesHelper(int magnitude)
+{
+ thrust::constant_iterator<long long> begin(1);
+ thrust::constant_iterator<long long> end = begin + (1ll << magnitude);
+ ASSERT_EQUAL(thrust::distance(begin, end), 1ll << magnitude);
+
+ long long result = thrust::reduce(thrust::device, begin, end);
+
+ ASSERT_EQUAL(result, 1ll << magnitude);
+}
+
+void TestReduceWithBigIndexes()
+{
+ TestReduceWithBigIndexesHelper(30);
+ TestReduceWithBigIndexesHelper(31);
+ TestReduceWithBigIndexesHelper(32);
+ TestReduceWithBigIndexesHelper(33);
+}
+DECLARE_UNITTEST(TestReduceWithBigIndexes);

The great Thrust index type fix, part 7: partition point, reduce tests.