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 403862256..27f8b53bd 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) 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..f7004f8c7 --- /dev/null +++ b/testing/transform_output_iterator_reduce_by_key.cu @@ -0,0 +1,51 @@ +#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; + diff --git a/thrust/iterator/transform_input_output_iterator.h b/thrust/iterator/transform_input_output_iterator.h index f512a36cb..a5f725dc5 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 * diff --git a/thrust/iterator/transform_output_iterator.h b/thrust/iterator/transform_output_iterator.h index 66fb46a37..3ac4b8572 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 *