The great Thrust index type fix, part 5: verify equal, find, IP.

Author: griwes

testing/equal.cu

@@ -2,6 +2,8 @@
 #include <thrust/equal.h>
 #include <thrust/functional.h>
 #include <thrust/iterator/retag.h>
+#include <thrust/device_malloc.h>
+#include <thrust/device_free.h>
 
 template <class Vector>
 void TestEqualSimple(void)
@@ -102,3 +104,48 @@ void TestEqualDispatchImplicit()
 }
 DECLARE_UNITTEST(TestEqualDispatchImplicit);
 
+struct only_set_when_both_expected
+{
+    long long expected;
+    bool * flag;
+
+    __device__
+    bool operator()(long long x, long long y)
+    {
+        if (x == expected && y == expected)
+        {
+            *flag = true;
+        }
+
+        return x == y;
+    }
+};
+
+void TestEqualWithBigIndexesHelper(int magnitude)
+{
+    thrust::counting_iterator<long long> begin(1);
+    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_both_expected fn = { (1ll << magnitude) - 1,
+        thrust::raw_pointer_cast(has_executed) };
+
+    ASSERT_EQUAL(thrust::equal(thrust::device, begin, end, begin, fn), true);
+
+    bool has_executed_h = *has_executed;
+    thrust::device_free(has_executed);
+
+    ASSERT_EQUAL(has_executed_h, true);
+}
+
+void TestEqualWithBigIndexes()
+{
+    TestEqualWithBigIndexesHelper(30);
+    TestEqualWithBigIndexesHelper(31);
+    TestEqualWithBigIndexesHelper(32);
+    TestEqualWithBigIndexesHelper(33);
+}
+DECLARE_UNITTEST(TestEqualWithBigIndexes);

testing/find.cu

@@ -304,3 +304,37 @@ struct TestFindIfNot
 };
 VariableUnitTest<TestFindIfNot, SignedIntegralTypes> TestFindIfNotInstance;
 
+void TestFindWithBigIndexesHelper(int magnitude)
+{
+    thrust::counting_iterator<long long> begin(1);
+    thrust::counting_iterator<long long> end = begin + (1ll << magnitude);
+    ASSERT_EQUAL(thrust::distance(begin, end), 1ll << magnitude);
+
+    thrust::detail::intmax_t distance_low_value = thrust::distance(
+        begin,
+        thrust::find(
+            thrust::device,
+            begin,
+            end,
+            17));
+
+    thrust::detail::intmax_t distance_high_value = thrust::distance(
+        begin,
+        thrust::find(
+            thrust::device,
+            begin,
+            end,
+            (1ll << magnitude) - 17));
+
+    ASSERT_EQUAL(distance_low_value, 16);
+    ASSERT_EQUAL(distance_high_value, (1ll << magnitude) - 18);
+}
+
+void TestFindWithBigIndexes()
+{
+    TestFindWithBigIndexesHelper(30);
+    TestFindWithBigIndexesHelper(31);
+    TestFindWithBigIndexesHelper(32);
+    TestFindWithBigIndexesHelper(33);
+}
+DECLARE_UNITTEST(TestFindWithBigIndexes);

testing/inner_product.cu

@@ -1,6 +1,8 @@
 #include <unittest/unittest.h>
 #include <thrust/inner_product.h>
 #include <thrust/iterator/retag.h>
+#include <thrust/device_malloc.h>
+#include <thrust/device_free.h>
 
 template <class Vector>
 void TestInnerProductSimple(void)
@@ -100,4 +102,54 @@ struct TestInnerProduct
 };
 VariableUnitTest<TestInnerProduct, IntegralTypes> TestInnerProductInstance;
 
+struct only_set_when_both_expected
+{
+    long long expected;
+    bool * flag;
+
+    __device__
+    long long operator()(long long x, long long y)
+    {
+        if (x == expected && y == expected)
+        {
+            *flag = true;
+        }
+
+        return x == y;
+    }
+};
+
+void TestInnerProductWithBigIndexesHelper(int magnitude)
+{
+    thrust::counting_iterator<long long> begin(1);
+    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_both_expected fn = { (1ll << magnitude) - 1,
+        thrust::raw_pointer_cast(has_executed) };
+
+    ASSERT_EQUAL(thrust::inner_product(
+        thrust::device,
+        begin, end,
+        begin,
+        0ll,
+        thrust::plus<long long>(),
+        fn), (1ll << magnitude));
+
+    bool has_executed_h = *has_executed;
+    thrust::device_free(has_executed);
+
+    ASSERT_EQUAL(has_executed_h, true);
+}
+
+void TestInnerProductWithBigIndexes()
+{
+    TestInnerProductWithBigIndexesHelper(30);
+    TestInnerProductWithBigIndexesHelper(31);
+    TestInnerProductWithBigIndexesHelper(32);
+    TestInnerProductWithBigIndexesHelper(33);
+}
+DECLARE_UNITTEST(TestInnerProductWithBigIndexes);