#include #include #include #include #include #include #include template OutputIterator transform_inclusive_scan(my_system &system, InputIterator, InputIterator, OutputIterator result, UnaryFunction, AssociativeOperator) { system.validate_dispatch(); return result; } void TestTransformInclusiveScanDispatchExplicit() { thrust::device_vector vec(1); my_system sys(0); thrust::transform_inclusive_scan(sys, vec.begin(), vec.begin(), vec.begin(), 0, 0); ASSERT_EQUAL(true, sys.is_valid()); } DECLARE_UNITTEST(TestTransformInclusiveScanDispatchExplicit); template OutputIterator transform_inclusive_scan(my_tag, InputIterator, InputIterator, OutputIterator result, UnaryFunction, AssociativeOperator) { *result = 13; return result; } void TestTransformInclusiveScanDispatchImplicit() { thrust::device_vector vec(1); thrust::transform_inclusive_scan(thrust::retag(vec.begin()), thrust::retag(vec.begin()), thrust::retag(vec.begin()), 0, 0); ASSERT_EQUAL(13, vec.front()); } DECLARE_UNITTEST(TestTransformInclusiveScanDispatchImplicit); template OutputIterator transform_exclusive_scan(my_system &system, InputIterator, InputIterator, OutputIterator result, UnaryFunction, T, AssociativeOperator) { system.validate_dispatch(); return result; } void TestTransformExclusiveScanDispatchExplicit() { thrust::device_vector vec(1); my_system sys(0); thrust::transform_exclusive_scan(sys, vec.begin(), vec.begin(), vec.begin(), 0, 0, 0); ASSERT_EQUAL(true, sys.is_valid()); } DECLARE_UNITTEST(TestTransformExclusiveScanDispatchExplicit); template OutputIterator transform_exclusive_scan(my_tag, InputIterator, InputIterator, OutputIterator result, UnaryFunction, T, AssociativeOperator) { *result = 13; return result; } void TestTransformExclusiveScanDispatchImplicit() { thrust::device_vector vec(1); thrust::transform_exclusive_scan(thrust::retag(vec.begin()), thrust::retag(vec.begin()), thrust::retag(vec.begin()), 0, 0, 0); ASSERT_EQUAL(13, vec.front()); } DECLARE_UNITTEST(TestTransformExclusiveScanDispatchImplicit); template void TestTransformScanSimple(void) { typedef typename Vector::value_type T; typename Vector::iterator iter; Vector input(5); Vector result(5); Vector output(5); input[0] = 1; input[1] = 3; input[2] = -2; input[3] = 4; input[4] = -5; Vector input_copy(input); // inclusive scan iter = thrust::transform_inclusive_scan(input.begin(), input.end(), output.begin(), thrust::negate(), thrust::plus()); result[0] = -1; result[1] = -4; result[2] = -2; result[3] = -6; result[4] = -1; ASSERT_EQUAL(std::size_t(iter - output.begin()), input.size()); ASSERT_EQUAL(input, input_copy); ASSERT_EQUAL(output, result); // exclusive scan with 0 init iter = thrust::transform_exclusive_scan(input.begin(), input.end(), output.begin(), thrust::negate(), 0, thrust::plus()); result[0] = 0; result[1] = -1; result[2] = -4; result[3] = -2; result[4] = -6; ASSERT_EQUAL(std::size_t(iter - output.begin()), input.size()); ASSERT_EQUAL(input, input_copy); ASSERT_EQUAL(output, result); // exclusive scan with nonzero init iter = thrust::transform_exclusive_scan(input.begin(), input.end(), output.begin(), thrust::negate(), 3, thrust::plus()); result[0] = 3; result[1] = 2; result[2] = -1; result[3] = 1; result[4] = -3; ASSERT_EQUAL(std::size_t(iter - output.begin()), input.size()); ASSERT_EQUAL(input, input_copy); ASSERT_EQUAL(output, result); // inplace inclusive scan input = input_copy; iter = thrust::transform_inclusive_scan(input.begin(), input.end(), input.begin(), thrust::negate(), thrust::plus()); result[0] = -1; result[1] = -4; result[2] = -2; result[3] = -6; result[4] = -1; ASSERT_EQUAL(std::size_t(iter - input.begin()), input.size()); ASSERT_EQUAL(input, result); // inplace exclusive scan with init input = input_copy; iter = thrust::transform_exclusive_scan(input.begin(), input.end(), input.begin(), thrust::negate(), 3, thrust::plus()); result[0] = 3; result[1] = 2; result[2] = -1; result[3] = 1; result[4] = -3; ASSERT_EQUAL(std::size_t(iter - input.begin()), input.size()); ASSERT_EQUAL(input, result); } DECLARE_INTEGRAL_VECTOR_UNITTEST(TestTransformScanSimple); struct Record { int number; bool operator==(const Record& rhs) const { return number == rhs.number; } bool operator!=(const Record& rhs) const { return !(rhs == *this); } friend Record operator+(Record lhs, const Record& rhs) { lhs.number += rhs.number; return lhs; } friend std::ostream& operator<<(std::ostream& os, const Record& record) { os << "number: " << record.number; return os; } }; struct negate { __host__ __device__ int operator()(Record const& record) const { return - record.number; } }; void TestTransformInclusiveScanDifferentTypes() { typename thrust::host_vector::iterator h_iter; thrust::host_vector h_input(5); thrust::host_vector h_output(5); thrust::host_vector result(5); h_input[0] = {1}; h_input[1] = {3}; h_input[2] = {-2}; h_input[3] = {4}; h_input[4] = {-5}; thrust::host_vector input_copy(h_input); h_iter = thrust::transform_inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), negate{}, thrust::plus{}); result[0] = -1; result[1] = -4; result[2] = -2; result[3] = -6; result[4] = -1; ASSERT_EQUAL(std::size_t(h_iter - h_output.begin()), h_input.size()); ASSERT_EQUAL(h_input, input_copy); ASSERT_EQUAL(h_output, result); typename thrust::device_vector::iterator d_iter; thrust::device_vector d_input = h_input; thrust::device_vector d_output(5); d_iter = thrust::transform_inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), negate{}, thrust::plus{}); ASSERT_EQUAL(std::size_t(d_iter - d_output.begin()), d_input.size()); ASSERT_EQUAL(d_input, input_copy); ASSERT_EQUAL(d_output, result); } DECLARE_UNITTEST(TestTransformInclusiveScanDifferentTypes); template struct TestTransformScan { void operator()(const size_t n) { thrust::host_vector h_input = unittest::random_integers(n); thrust::device_vector d_input = h_input; thrust::host_vector h_output(n); thrust::device_vector d_output(n); thrust::transform_inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), thrust::negate(), thrust::plus()); thrust::transform_inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), thrust::negate(), thrust::plus()); ASSERT_EQUAL(d_output, h_output); thrust::transform_exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), thrust::negate(), (T) 11, thrust::plus()); thrust::transform_exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), thrust::negate(), (T) 11, thrust::plus()); ASSERT_EQUAL(d_output, h_output); // in-place scans h_output = h_input; d_output = d_input; thrust::transform_inclusive_scan(h_output.begin(), h_output.end(), h_output.begin(), thrust::negate(), thrust::plus()); thrust::transform_inclusive_scan(d_output.begin(), d_output.end(), d_output.begin(), thrust::negate(), thrust::plus()); ASSERT_EQUAL(d_output, h_output); h_output = h_input; d_output = d_input; thrust::transform_exclusive_scan(h_output.begin(), h_output.end(), h_output.begin(), thrust::negate(), (T) 11, thrust::plus()); thrust::transform_exclusive_scan(d_output.begin(), d_output.end(), d_output.begin(), thrust::negate(), (T) 11, thrust::plus()); ASSERT_EQUAL(d_output, h_output); } }; VariableUnitTest TestTransformScanInstance; template void TestTransformScanCountingIterator(void) { typedef typename Vector::value_type T; typedef typename thrust::iterator_system::type space; thrust::counting_iterator first(1); Vector result(3); thrust::transform_inclusive_scan(first, first + 3, result.begin(), thrust::negate(), thrust::plus()); ASSERT_EQUAL(result[0], -1); ASSERT_EQUAL(result[1], -3); ASSERT_EQUAL(result[2], -6); } DECLARE_INTEGRAL_VECTOR_UNITTEST(TestTransformScanCountingIterator); template struct TestTransformScanToDiscardIterator { void operator()(const size_t n) { thrust::host_vector h_input = unittest::random_integers(n); thrust::device_vector d_input = h_input; thrust::discard_iterator<> reference(n); thrust::discard_iterator<> h_result = thrust::transform_inclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), thrust::negate(), thrust::plus()); thrust::discard_iterator<> d_result = thrust::transform_inclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), thrust::negate(), thrust::plus()); ASSERT_EQUAL_QUIET(reference, h_result); ASSERT_EQUAL_QUIET(reference, d_result); h_result = thrust::transform_exclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), thrust::negate(), (T) 11, thrust::plus()); d_result = thrust::transform_exclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), thrust::negate(), (T) 11, thrust::plus()); ASSERT_EQUAL_QUIET(reference, h_result); ASSERT_EQUAL_QUIET(reference, d_result); } }; VariableUnitTest TestTransformScanToDiscardIteratorInstance;