From 6cdb69d5d49f97b1ca9426d24891b183e8f43a28 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Oct 2022 17:12:12 +1100 Subject: [PATCH 1/7] Add default ctor to transform[_input]_output_iterator and add test --- testing/transform_output_iterator.cu | 168 +++++++++++------- .../transform_input_output_iterator.h | 75 ++++---- thrust/iterator/transform_output_iterator.h | 74 ++++---- 3 files changed, 181 insertions(+), 136 deletions(-) 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 - From daf61336f93968c7aff2687a3562f80c2ce24f69 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Oct 2022 17:50:53 +1100 Subject: [PATCH 2/7] Use =default --- thrust/iterator/transform_input_output_iterator.h | 4 +--- thrust/iterator/transform_output_iterator.h | 4 +--- 2 files changed, 2 insertions(+), 6 deletions(-) diff --git a/thrust/iterator/transform_input_output_iterator.h b/thrust/iterator/transform_input_output_iterator.h index 1a727feda..e6ae99c8c 100644 --- a/thrust/iterator/transform_input_output_iterator.h +++ b/thrust/iterator/transform_input_output_iterator.h @@ -100,9 +100,7 @@ class transform_input_output_iterator /*! \endcond */ - /*! Null constructor does nothing. - */ - __host__ __device__ transform_input_output_iterator() {} + transform_input_output_iterator() = default; /*! 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 diff --git a/thrust/iterator/transform_output_iterator.h b/thrust/iterator/transform_output_iterator.h index 791ba5eec..9c644ac02 100644 --- a/thrust/iterator/transform_output_iterator.h +++ b/thrust/iterator/transform_output_iterator.h @@ -102,9 +102,7 @@ class transform_output_iterator /*! \endcond */ - /*! Null constructor does nothing. - */ - __host__ __device__ transform_output_iterator() {} + transform_output_iterator() = default; /*! This constructor takes as argument an \c OutputIterator and an \c * UnaryFunction and copies them to a new \p transform_output_iterator From cd37987f63e88d17f7f095bd5b40012a889d707c Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 11 Oct 2022 12:05:45 +1100 Subject: [PATCH 3/7] Revert "Use =default" This reverts commit daf61336f93968c7aff2687a3562f80c2ce24f69. --- thrust/iterator/transform_input_output_iterator.h | 4 +++- thrust/iterator/transform_output_iterator.h | 4 +++- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/thrust/iterator/transform_input_output_iterator.h b/thrust/iterator/transform_input_output_iterator.h index e6ae99c8c..1a727feda 100644 --- a/thrust/iterator/transform_input_output_iterator.h +++ b/thrust/iterator/transform_input_output_iterator.h @@ -100,7 +100,9 @@ class transform_input_output_iterator /*! \endcond */ - transform_input_output_iterator() = default; + /*! 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 diff --git a/thrust/iterator/transform_output_iterator.h b/thrust/iterator/transform_output_iterator.h index 9c644ac02..791ba5eec 100644 --- a/thrust/iterator/transform_output_iterator.h +++ b/thrust/iterator/transform_output_iterator.h @@ -102,7 +102,9 @@ class transform_output_iterator /*! \endcond */ - transform_output_iterator() = default; + /*! 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 From bb9c6627deca02a228b92c95b49ed018d7a458fc Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 11 Oct 2022 12:05:57 +1100 Subject: [PATCH 4/7] Revert "Add default ctor to transform[_input]_output_iterator and add test" This reverts commit 6cdb69d5d49f97b1ca9426d24891b183e8f43a28. --- testing/transform_output_iterator.cu | 168 +++++++----------- .../transform_input_output_iterator.h | 75 ++++---- thrust/iterator/transform_output_iterator.h | 74 ++++---- 3 files changed, 136 insertions(+), 181 deletions(-) diff --git a/testing/transform_output_iterator.cu b/testing/transform_output_iterator.cu index e001278dc..403862256 100644 --- a/testing/transform_output_iterator.cu +++ b/testing/transform_output_iterator.cu @@ -1,133 +1,91 @@ -#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 thrust::square UnaryFunction; - typedef typename Vector::iterator Iterator; + typedef typename Vector::value_type T; - Vector input(4); - Vector output(4); + typedef thrust::square UnaryFunction; + typedef typename Vector::iterator Iterator; - // initialize input - thrust::sequence(input.begin(), input.end(), T{1}); + 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()); - // construct transform_iterator - thrust::transform_output_iterator output_iter(output.begin(), - UnaryFunction()); + thrust::copy(input.begin(), input.end(), output_iter); - thrust::copy(input.begin(), input.end(), output_iter); + Vector gold_output(4); + gold_output[0] = 1; + gold_output[1] = 4; + gold_output[2] = 9; + gold_output[3] = 16; - 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); - 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); + 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); - 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 1a727feda..f512a36cb 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,24 +85,22 @@ 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: - friend class thrust::iterator_core_access; - /*! \endcond - */ + typedef typename + detail::transform_input_output_iterator_base::type + super_t; - /*! Null constructor does nothing. + friend class thrust::iterator_core_access; + /*! \endcond */ - __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 @@ -112,30 +110,29 @@ class transform_input_output_iterator * \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( - 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< + InputFunction, OutputFunction, Iterator + >(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 @@ -149,13 +146,10 @@ class transform_input_output_iterator */ 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 @@ -165,3 +159,4 @@ transform_input_output_iterator */ THRUST_NAMESPACE_END + diff --git a/thrust/iterator/transform_output_iterator.h b/thrust/iterator/transform_output_iterator.h index 791ba5eec..66fb46a37 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: - friend class thrust::iterator_core_access; - /*! \endcond - */ + typedef typename + detail::transform_output_iterator_base::type + super_t; - /*! Null constructor does nothing. + friend class thrust::iterator_core_access; + /*! \endcond */ - __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( - 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< + UnaryFunction, OutputIterator + >(this->base_reference(), fun); + } + + UnaryFunction fun; + + /*! \endcond + */ }; // end transform_output_iterator /*! \p make_transform_output_iterator creates a \p transform_output_iterator from @@ -146,9 +146,10 @@ class transform_output_iterator */ 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 @@ -158,3 +159,4 @@ transform_output_iterator */ THRUST_NAMESPACE_END + From 65aeefe67d279c141222057ab9ccbe849827ac83 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 11 Oct 2022 12:11:51 +1100 Subject: [PATCH 5/7] Add default ctors and test --- testing/transform_output_iterator.cu | 36 +++++++++++++++++++ .../transform_input_output_iterator.h | 3 +- thrust/iterator/transform_output_iterator.h | 3 +- 3 files changed, 40 insertions(+), 2 deletions(-) diff --git a/testing/transform_output_iterator.cu b/testing/transform_output_iterator.cu index 403862256..1cf21a201 100644 --- a/testing/transform_output_iterator.cu +++ b/testing/transform_output_iterator.cu @@ -89,3 +89,39 @@ struct TestTransformOutputIteratorScan }; 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..8db86f6f0 100644 --- a/thrust/iterator/transform_input_output_iterator.h +++ b/thrust/iterator/transform_input_output_iterator.h @@ -102,6 +102,8 @@ template /*! \endcond */ + transform_input_output_iterator() = default; + /*! 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 * @@ -159,4 +161,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..e6111a409 100644 --- a/thrust/iterator/transform_output_iterator.h +++ b/thrust/iterator/transform_output_iterator.h @@ -104,6 +104,8 @@ template /*! \endcond */ + transform_output_iterator() = default; + /*! This constructor takes as argument an \c OutputIterator and an \c * UnaryFunction and copies them to a new \p transform_output_iterator * @@ -159,4 +161,3 @@ make_transform_output_iterator(OutputIterator out, UnaryFunction fun) */ THRUST_NAMESPACE_END - From 03e988cfe64e3edc064b9c5083ba83b16dec97b6 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 11 Oct 2022 13:25:00 +1100 Subject: [PATCH 6/7] Split new test into its own file and disable on TBB for now. --- testing/CMakeLists.txt | 4 ++ testing/transform_output_iterator.cu | 45 ++--------------- ...transform_output_iterator_reduce_by_key.cu | 50 +++++++++++++++++++ 3 files changed, 59 insertions(+), 40 deletions(-) create mode 100644 testing/transform_output_iterator_reduce_by_key.cu diff --git a/testing/CMakeLists.txt b/testing/CMakeLists.txt index af60c5442..69f870d8c 100644 --- a/testing/CMakeLists.txt +++ b/testing/CMakeLists.txt @@ -41,6 +41,10 @@ thrust_declare_test_restrictions(future CPP.CUDA OMP.CUDA TBB.CUDA) # for CUDA. thrust_declare_test_restrictions(unittest_static_assert CPP.CPP CPP.CUDA) +# In the TBB backend, reduce_by_key does not currently work with transform_output_iterator +# https://github.com/NVIDIA/thrust/issues/1811 +thrust_declare_test_restrictions(transform_output_iterator_reduce_by_key CPP.CPP CPP.OMP CPP.CUDA) + ## thrust_add_test # # Add a test executable and register it with ctest. diff --git a/testing/transform_output_iterator.cu b/testing/transform_output_iterator.cu index 1cf21a201..3de12e155 100644 --- a/testing/transform_output_iterator.cu +++ b/testing/transform_output_iterator.cu @@ -1,11 +1,13 @@ #include -#include #include -#include +#include #include -#include +#include #include +#include +#include +#include template void TestTransformOutputIterator(void) @@ -88,40 +90,3 @@ struct TestTransformOutputIteratorScan } }; 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/testing/transform_output_iterator_reduce_by_key.cu b/testing/transform_output_iterator_reduce_by_key.cu new file mode 100644 index 000000000..7018d76d0 --- /dev/null +++ b/testing/transform_output_iterator_reduce_by_key.cu @@ -0,0 +1,50 @@ +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + + +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; From abd0bed6f5c7adf1a0873e32710b2b77997b3abd Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 11 Oct 2022 17:25:22 +1100 Subject: [PATCH 7/7] re-add redundant newlines at eof --- testing/transform_output_iterator.cu | 1 + testing/transform_output_iterator_reduce_by_key.cu | 1 + thrust/iterator/transform_input_output_iterator.h | 1 + thrust/iterator/transform_output_iterator.h | 1 + 4 files changed, 4 insertions(+) diff --git a/testing/transform_output_iterator.cu b/testing/transform_output_iterator.cu index 3de12e155..27f8b53bd 100644 --- a/testing/transform_output_iterator.cu +++ b/testing/transform_output_iterator.cu @@ -90,3 +90,4 @@ struct TestTransformOutputIteratorScan } }; VariableUnitTest TestTransformOutputIteratorScanInstance; + diff --git a/testing/transform_output_iterator_reduce_by_key.cu b/testing/transform_output_iterator_reduce_by_key.cu index 7018d76d0..f7004f8c7 100644 --- a/testing/transform_output_iterator_reduce_by_key.cu +++ b/testing/transform_output_iterator_reduce_by_key.cu @@ -48,3 +48,4 @@ struct TestTransformOutputIteratorReduceByKey }; VariableUnitTest TestTransformOutputIteratorReduceByKeyInstance; + diff --git a/thrust/iterator/transform_input_output_iterator.h b/thrust/iterator/transform_input_output_iterator.h index 8db86f6f0..a5f725dc5 100644 --- a/thrust/iterator/transform_input_output_iterator.h +++ b/thrust/iterator/transform_input_output_iterator.h @@ -161,3 +161,4 @@ 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 e6111a409..3ac4b8572 100644 --- a/thrust/iterator/transform_output_iterator.h +++ b/thrust/iterator/transform_output_iterator.h @@ -161,3 +161,4 @@ make_transform_output_iterator(OutputIterator out, UnaryFunction fun) */ THRUST_NAMESPACE_END +