The great Thrust index type fix, part 3: (partially) verify copy.
MODIFY
testing/copy.cu
testing/copy.cu
#include <thrust/iterator/constant_iterator.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/retag.h>
+#include <thrust/device_malloc.h>
+#include <thrust/device_free.h>
void TestCopyFromConstIterator(void)
{
@@ -648,3 +650,63 @@ void TestCopyIfStencilDispatchImplicit()
ASSERT_EQUAL(13, vec.front());
}
DECLARE_UNITTEST(TestCopyIfStencilDispatchImplicit);
+
+struct only_set_when_expected_it
+{
+ unsigned long long expected;
+ bool * flag;
+
+ __host__ __device__ only_set_when_expected_it operator++() const { return *this; }
+ __host__ __device__ only_set_when_expected_it operator*() const { return *this; }
+ template<typename Difference>
+ __host__ __device__ only_set_when_expected_it operator+(Difference) const { return *this; }
+ template<typename Index>
+ __host__ __device__ only_set_when_expected_it operator[](Index) const { return *this; }
+
+ __device__
+ void operator=(long long value) const
+ {
+ if (value == expected)
+ {
+ *flag = true;
+ }
+ }
+};
+
+namespace thrust
+{
+template<>
+struct iterator_traits<only_set_when_expected_it>
+{
+ typedef long long value_type;
+ typedef only_set_when_expected_it reference;
+};
+}
+
+void TestCopyWithBigIndexesHelper(int magnitude)
+{
+ thrust::counting_iterator<unsigned long long> begin(0);
+ thrust::counting_iterator<unsigned long long> end = begin + (1ull << 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) };
+
+ thrust::copy(thrust::device, begin, end, out);
+
+ bool has_executed_h = *has_executed;
+ thrust::device_free(has_executed);
+
+ ASSERT_EQUAL(has_executed_h, true);
+}
+
+void TestCopyWithBigIndexes()
+{
+ TestCopyWithBigIndexesHelper(30);
+ TestCopyWithBigIndexesHelper(31);
+ TestCopyWithBigIndexesHelper(32);
+ TestCopyWithBigIndexesHelper(33);
+}
+DECLARE_UNITTEST(TestCopyWithBigIndexes);

MODIFY
testing/for_each.cu
testing/for_each.cu
@@ -372,7 +372,7 @@ DECLARE_UNITTEST(TestForEachNWithLargeTypes);
THRUST_DISABLE_MSVC_POSSIBLE_LOSS_OF_DATA_WARNING_END
-struct OnlySetWhenExpected
+struct only_set_when_expected
{
unsigned long long expected;
bool * flag;
@@ -396,7 +396,7 @@ void TestForEachWithBigIndexesHelper(int magnitude)
thrust::device_ptr<bool> has_executed = thrust::device_malloc<bool>(1);
*has_executed = false;
- OnlySetWhenExpected fn = { (1ull << magnitude) - 1, thrust::raw_pointer_cast(has_executed) };
+ only_set_when_expected fn = { (1ull << magnitude) - 1, thrust::raw_pointer_cast(has_executed) };
thrust::for_each(thrust::device, begin, end, fn);

The great Thrust index type fix, part 3: (partially) verify copy.