diff --git a/testing/transform_output_iterator.cu b/testing/transform_output_iterator.cu index 403862256..e001278dc 100644 --- a/testing/transform_output_iterator.cu +++ b/testing/transform_output_iterator.cu @@ -1,91 +1,133 @@ -#include -#include - #include -#include +#include #include -#include +#include #include +#include +#include +#include +#include +#include + +#include +#include template void TestTransformOutputIterator(void) { - typedef typename Vector::value_type T; + typedef typename Vector::value_type T; + + typedef thrust::square UnaryFunction; + typedef typename Vector::iterator Iterator; - typedef thrust::square UnaryFunction; - typedef typename Vector::iterator Iterator; + Vector input(4); + Vector output(4); - Vector input(4); - Vector output(4); - - // initialize input - thrust::sequence(input.begin(), input.end(), T{1}); - - // construct transform_iterator - thrust::transform_output_iterator output_iter(output.begin(), UnaryFunction()); + // initialize input + thrust::sequence(input.begin(), input.end(), T{1}); - thrust::copy(input.begin(), input.end(), output_iter); + // construct transform_iterator + thrust::transform_output_iterator output_iter(output.begin(), + UnaryFunction()); - Vector gold_output(4); - gold_output[0] = 1; - gold_output[1] = 4; - gold_output[2] = 9; - gold_output[3] = 16; + thrust::copy(input.begin(), input.end(), output_iter); - ASSERT_EQUAL(output, gold_output); + Vector gold_output(4); + gold_output[0] = 1; + gold_output[1] = 4; + gold_output[2] = 9; + gold_output[3] = 16; + ASSERT_EQUAL(output, gold_output); } DECLARE_VECTOR_UNITTEST(TestTransformOutputIterator); template void TestMakeTransformOutputIterator(void) { - typedef typename Vector::value_type T; - - typedef thrust::square UnaryFunction; - - Vector input(4); - Vector output(4); - - // initialize input - thrust::sequence(input.begin(), input.end(), 1); - - thrust::copy(input.begin(), input.end(), - thrust::make_transform_output_iterator(output.begin(), UnaryFunction())); - - Vector gold_output(4); - gold_output[0] = 1; - gold_output[1] = 4; - gold_output[2] = 9; - gold_output[3] = 16; - ASSERT_EQUAL(output, gold_output); + typedef typename Vector::value_type T; + + typedef thrust::square UnaryFunction; + + Vector input(4); + Vector output(4); + + // initialize input + thrust::sequence(input.begin(), input.end(), 1); + thrust::copy(input.begin(), + input.end(), + thrust::make_transform_output_iterator(output.begin(), UnaryFunction())); + + Vector gold_output(4); + gold_output[0] = 1; + gold_output[1] = 4; + gold_output[2] = 9; + gold_output[3] = 16; + ASSERT_EQUAL(output, gold_output); } DECLARE_VECTOR_UNITTEST(TestMakeTransformOutputIterator); template struct TestTransformOutputIteratorScan { - void operator()(const size_t n) - { - thrust::host_vector h_data = unittest::random_samples(n); - thrust::device_vector d_data = h_data; - - thrust::host_vector h_result(n); - thrust::device_vector d_result(n); - - // run on host - thrust::inclusive_scan(thrust::make_transform_iterator(h_data.begin(), thrust::negate()), - thrust::make_transform_iterator(h_data.end(), thrust::negate()), - h_result.begin()); - // run on device - thrust::inclusive_scan(d_data.begin(), d_data.end(), - thrust::make_transform_output_iterator( - d_result.begin(), thrust::negate())); - - - ASSERT_EQUAL(h_result, d_result); - } + void operator()(const size_t n) + { + thrust::host_vector h_data = unittest::random_samples(n); + thrust::device_vector d_data = h_data; + + thrust::host_vector h_result(n); + thrust::device_vector d_result(n); + + // run on host + thrust::inclusive_scan(thrust::make_transform_iterator(h_data.begin(), thrust::negate()), + thrust::make_transform_iterator(h_data.end(), thrust::negate()), + h_result.begin()); + // run on device + thrust::inclusive_scan(d_data.begin(), + d_data.end(), + thrust::make_transform_output_iterator(d_result.begin(), + thrust::negate())); + + ASSERT_EQUAL(h_result, d_result); + } }; -VariableUnitTest TestTransformOutputIteratorScanInstance; +VariableUnitTest + TestTransformOutputIteratorScanInstance; +template +struct TestTransformOutputIteratorReduceByKey +{ + void operator()(const size_t n) + { + thrust::host_vector h_keys = unittest::random_samples(n); + thrust::sort(h_keys.begin(), h_keys.end()); + thrust::device_vector d_keys = h_keys; + + thrust::host_vector h_values = unittest::random_samples(n); + thrust::device_vector d_values = h_values; + + thrust::host_vector h_result(n); + thrust::device_vector d_result(n); + + // run on host + thrust::reduce_by_key(thrust::host, + h_keys.begin(), + h_keys.end(), + thrust::make_transform_iterator(h_values.begin(), thrust::negate()), + thrust::discard_iterator{}, + h_result.begin()); + // run on device + thrust::reduce_by_key(thrust::device, + d_keys.begin(), + d_keys.end(), + d_values.begin(), + thrust::discard_iterator{}, + thrust::make_transform_output_iterator(d_result.begin(), + thrust::negate())); + + ASSERT_EQUAL(h_result, d_result); + } +}; +VariableUnitTest + TestTransformOutputIteratorReduceByKeyInstance; diff --git a/thrust/iterator/transform_input_output_iterator.h b/thrust/iterator/transform_input_output_iterator.h index f512a36cb..1a727feda 100644 --- a/thrust/iterator/transform_input_output_iterator.h +++ b/thrust/iterator/transform_input_output_iterator.h @@ -62,7 +62,7 @@ THRUST_NAMESPACE_BEGIN * // Iterator that returns negated values and writes squared values * auto iter = thrust::make_transform_input_output_iterator(v.begin(), * thrust::negate{}, thrust::square{}); - * + * * // Iterator negates values when reading * std::cout << iter[0] << " "; // -1.0f; * std::cout << iter[1] << " "; // -2.0f; @@ -85,23 +85,25 @@ THRUST_NAMESPACE_BEGIN */ template - class transform_input_output_iterator +class transform_input_output_iterator : public detail::transform_input_output_iterator_base::type { /*! \cond */ - public: - - typedef typename - detail::transform_input_output_iterator_base::type - super_t; +public: + typedef typename detail:: + transform_input_output_iterator_base::type super_t; - friend class thrust::iterator_core_access; + friend class thrust::iterator_core_access; /*! \endcond */ + /*! Null constructor does nothing. + */ + __host__ __device__ transform_input_output_iterator() {} + /*! This constructor takes as argument a \c Iterator an \c InputFunction and an * \c OutputFunction and copies them to a new \p transform_input_output_iterator * @@ -110,29 +112,30 @@ template * \param input_function An \c InputFunction to be executed on values read from the iterator * \param output_function An \c OutputFunction to be executed on values written to the iterator */ - __host__ __device__ - transform_input_output_iterator(Iterator const& io, InputFunction input_function, OutputFunction output_function) - : super_t(io), input_function(input_function), output_function(output_function) - { - } - - /*! \cond - */ - private: - - __host__ __device__ - typename super_t::reference dereference() const - { - return detail::transform_input_output_iterator_proxy< - InputFunction, OutputFunction, Iterator - >(this->base_reference(), input_function, output_function); - } - - InputFunction input_function; - OutputFunction output_function; - - /*! \endcond - */ + __host__ __device__ transform_input_output_iterator(Iterator const &io, + InputFunction input_function, + OutputFunction output_function) + : super_t(io) + , input_function(input_function) + , output_function(output_function) + {} + + /*! \cond + */ +private: + __host__ __device__ typename super_t::reference dereference() const + { + return detail::transform_input_output_iterator_proxy( + this->base_reference(), + input_function, + output_function); + } + + InputFunction input_function; + OutputFunction output_function; + + /*! \endcond + */ }; // end transform_input_output_iterator /*! \p make_transform_input_output_iterator creates a \p transform_input_output_iterator from @@ -146,10 +149,13 @@ template */ template transform_input_output_iterator -__host__ __device__ -make_transform_input_output_iterator(Iterator io, InputFunction input_function, OutputFunction output_function) + __host__ __device__ make_transform_input_output_iterator(Iterator io, + InputFunction input_function, + OutputFunction output_function) { - return transform_input_output_iterator(io, input_function, output_function); + return transform_input_output_iterator(io, + input_function, + output_function); } // end make_transform_input_output_iterator /*! \} // end fancyiterators @@ -159,4 +165,3 @@ make_transform_input_output_iterator(Iterator io, InputFunction input_function, */ THRUST_NAMESPACE_END - diff --git a/thrust/iterator/transform_output_iterator.h b/thrust/iterator/transform_output_iterator.h index 66fb46a37..791ba5eec 100644 --- a/thrust/iterator/transform_output_iterator.h +++ b/thrust/iterator/transform_output_iterator.h @@ -38,7 +38,7 @@ THRUST_NAMESPACE_BEGIN /*! \p transform_output_iterator is a special kind of output iterator which * transforms a value written upon dereference. This iterator is useful * for transforming an output from algorithms without explicitly storing the - * intermediate result in the memory and applying subsequent transformation, + * intermediate result in the memory and applying subsequent transformation, * thereby avoiding wasting memory capacity and bandwidth. * Using \p transform_iterator facilitates kernel fusion by deferring execution * of transformation until the value is written while saving both memory @@ -61,7 +61,7 @@ THRUST_NAMESPACE_BEGIN * return sqrtf(x); * } * }; - * + * * int main() * { * thrust::device_vector v(4); @@ -69,17 +69,17 @@ THRUST_NAMESPACE_BEGIN * typedef thrust::device_vector::iterator FloatIterator; * thrust::transform_output_iterator iter(v.begin(), square_root()); * - * iter[0] = 1.0f; // stores sqrtf( 1.0f) + * iter[0] = 1.0f; // stores sqrtf( 1.0f) * iter[1] = 4.0f; // stores sqrtf( 4.0f) * iter[2] = 9.0f; // stores sqrtf( 9.0f) * iter[3] = 16.0f; // stores sqrtf(16.0f) * // iter[4] is an out-of-bounds error - * + * * v[0]; // returns 1.0f; * v[1]; // returns 2.0f; * v[2]; // returns 3.0f; * v[3]; // returns 4.0f; - * + * * } * \endcode * @@ -87,52 +87,52 @@ THRUST_NAMESPACE_BEGIN */ template - class transform_output_iterator +class transform_output_iterator : public detail::transform_output_iterator_base::type { /*! \cond */ - public: - - typedef typename - detail::transform_output_iterator_base::type - super_t; +public: + typedef + typename detail::transform_output_iterator_base::type super_t; - friend class thrust::iterator_core_access; + friend class thrust::iterator_core_access; /*! \endcond */ + /*! Null constructor does nothing. + */ + __host__ __device__ transform_output_iterator() {} + /*! This constructor takes as argument an \c OutputIterator and an \c * UnaryFunction and copies them to a new \p transform_output_iterator * - * \param out An \c OutputIterator pointing to the output range whereto the result of + * \param out An \c OutputIterator pointing to the output range whereto the result of * \p transform_output_iterator's \c UnaryFunction will be written. * \param fun An \c UnaryFunction used to transform the objects assigned to * this \p transform_output_iterator. */ - __host__ __device__ - transform_output_iterator(OutputIterator const& out, UnaryFunction fun) : super_t(out), fun(fun) - { - } - - /*! \cond - */ - private: - - __host__ __device__ - typename super_t::reference dereference() const - { - return detail::transform_output_iterator_proxy< - UnaryFunction, OutputIterator - >(this->base_reference(), fun); - } - - UnaryFunction fun; - - /*! \endcond - */ + __host__ __device__ transform_output_iterator(OutputIterator const &out, UnaryFunction fun) + : super_t(out) + , fun(fun) + {} + + /*! \cond + */ +private: + __host__ __device__ typename super_t::reference dereference() const + { + return detail::transform_output_iterator_proxy( + this->base_reference(), + fun); + } + + UnaryFunction fun; + + /*! \endcond + */ }; // end transform_output_iterator /*! \p make_transform_output_iterator creates a \p transform_output_iterator from @@ -146,10 +146,9 @@ template */ template transform_output_iterator -__host__ __device__ -make_transform_output_iterator(OutputIterator out, UnaryFunction fun) + __host__ __device__ make_transform_output_iterator(OutputIterator out, UnaryFunction fun) { - return transform_output_iterator(out, fun); + return transform_output_iterator(out, fun); } // end make_transform_output_iterator /*! \} // end fancyiterators @@ -159,4 +158,3 @@ make_transform_output_iterator(OutputIterator out, UnaryFunction fun) */ THRUST_NAMESPACE_END -