Implement transform reduce and remove tests
Updated the execution policy header too.
Replaced the random sample in vector insert tests.
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -70,10 +70,12 @@ add_hipthrust_test("thrust.hip.for_each" test_for_each.cpp)
add_hipthrust_test("thrust.hip.memory" test_memory.cpp)
add_hipthrust_test("thrust.hip.mismatch" test_mismatch.cpp)
add_hipthrust_test("thrust.hip.reduce" test_reduce.cpp)
+add_hipthrust_test("thrust.hip.remove" test_remove.cpp)
add_hipthrust_test("thrust.hip.scatter" test_scatter.cpp)
add_hipthrust_test("thrust.hip.sort" test_sort.cpp)
add_hipthrust_test("thrust.hip.tabulate" test_tabulate.cpp)
add_hipthrust_test("thrust.hip.transform" test_transform.cpp)
+add_hipthrust_test("thrust.hip.transform_reduce" test_transform_reduce.cpp)
add_hipthrust_test("thrust.hip.tuple" test_tuple.cpp)
add_hipthrust_test("thrust.hip.uninitialized_copy" test_uninitialized_copy.cpp)
add_hipthrust_test("thrust.hip.uninitialized_fill" test_uninitialized_fill.cpp)

ADD
test/test_remove.cpp
+// MIT License
+//
+// Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#include <iostream>
+#include <type_traits>
+#include <cstdlib>
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/remove.h>
+
+#include <thrust/count.h>
+#include <thrust/functional.h>
+#include <stdexcept>
+#include <thrust/iterator/discard_iterator.h>
+#include <thrust/iterator/zip_iterator.h>
+#include <thrust/iterator/retag.h>
+
+#include "test_utils.hpp"
+
+template<
+ class Input
+>
+struct Params
+{
+ using input_type = Input;
+};
+
+template<class Params>
+class RemoveTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<thrust::host_vector<short>>,
+ Params<thrust::host_vector<int>>,
+ Params<thrust::host_vector<long long>>,
+ Params<thrust::host_vector<unsigned short>>,
+ Params<thrust::host_vector<unsigned int>>,
+ Params<thrust::host_vector<unsigned long long>>,
+ Params<thrust::host_vector<float>>,
+ Params<thrust::host_vector<double>>,
+ Params<thrust::device_vector<short>>,
+ Params<thrust::device_vector<int>>,
+ Params<thrust::device_vector<long long>>,
+ Params<thrust::device_vector<unsigned short>>,
+ Params<thrust::device_vector<unsigned int>>,
+ Params<thrust::device_vector<unsigned long long>>,
+ Params<thrust::device_vector<float>>,
+ Params<thrust::device_vector<double>>
+> RemoveTestsParams;
+
+TYPED_TEST_CASE(RemoveTests, RemoveTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+template<typename T>
+struct is_even
+ : thrust::unary_function<T,bool>
+{
+ __host__ __device__
+ bool operator()(T x) { return (static_cast<unsigned int>(x) & 1) == 0; }
+};
+
+template<typename T>
+struct is_true
+ : thrust::unary_function<T,bool>
+{
+ __host__ __device__
+ bool operator()(T x) { return x ? true : false; }
+};
+
+TYPED_TEST(RemoveTests, TestRemoveSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector data(5);
+ data[0] = T(1);
+ data[1] = T(2);
+ data[2] = T(1);
+ data[3] = T(3);
+ data[4] = T(2);
+
+ typename Vector::iterator end = thrust::remove(data.begin(),
+ data.end(),
+ (T) 2);
+
+ ASSERT_EQ(end - data.begin(), 3);
+
+ ASSERT_EQ(data[0], T(1));
+ ASSERT_EQ(data[1], T(1));
+ ASSERT_EQ(data[2], T(3));
+}
+
+template<typename ForwardIterator,
+ typename T>
+ForwardIterator remove(my_system &system,
+ ForwardIterator first,
+ ForwardIterator,
+ const T &)
+{
+ system.validate_dispatch();
+ return first;
+}
+
+TEST(RemoveTests, TestRemoveDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::remove(sys, vec.begin(), vec.end(), 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename ForwardIterator,
+ typename T>
+ForwardIterator remove(my_tag,
+ ForwardIterator first,
+ ForwardIterator,
+ const T &)
+{
+ *first = 13;
+ return first;
+}
+
+TEST(RemoveTests, TestRemoveDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::remove(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ 0);
+
+ ASSERT_EQ(13, vec.front());
+}
+
+TYPED_TEST(RemoveTests, TestRemoveCopySimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector data(5);
+ data[0] = T(1);
+ data[1] = T(2);
+ data[2] = T(1);
+ data[3] = T(3);
+ data[4] = T(2);
+
+ Vector result(5);
+
+ typename Vector::iterator end = thrust::remove_copy(data.begin(),
+ data.end(),
+ result.begin(),
+ T(2));
+
+ ASSERT_EQ(end - result.begin(), 3);
+
+ ASSERT_EQ(result[0], T(1));
+ ASSERT_EQ(result[1], T(1));
+ ASSERT_EQ(result[2], T(3));
+}
+
+template<typename InputIterator,
+ typename OutputIterator,
+ typename T>
+OutputIterator remove_copy(my_system &system,
+ InputIterator,
+ InputIterator,
+ OutputIterator result,
+ const T &)
+{
+ system.validate_dispatch();
+ return result;
+}
+
+TEST(RemoveTests, TestRemoveCopyDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::remove_copy(sys,
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename InputIterator,
+ typename OutputIterator,
+ typename T>
+OutputIterator remove_copy(my_tag,
+ InputIterator,
+ InputIterator,
+ OutputIterator result,
+ const T &)
+{
+ *result = 13;
+ return result;
+}
+
+TEST(RemoveTests, TestRemoveCopyDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::remove_copy(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ 0);
+
+ ASSERT_EQ(13, vec.front());
+}
+
+TYPED_TEST(RemoveTests, TestRemoveIfSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector data(5);
+ data[0] = 1;
+ data[1] = 2;
+ data[2] = 1;
+ data[3] = 3;
+ data[4] = 2;
+
+ typename Vector::iterator end = thrust::remove_if(data.begin(),
+ data.end(),
+ is_even<T>());
+
+ ASSERT_EQ(end - data.begin(), 3);
+
+ ASSERT_EQ(data[0], 1);
+ ASSERT_EQ(data[1], 1);
+ ASSERT_EQ(data[2], 3);
+}
+
+template<typename ForwardIterator,
+ typename Predicate>
+__host__ __device__
+ForwardIterator remove_if(my_system &system,
+ ForwardIterator first,
+ ForwardIterator,
+ Predicate)
+{
+ system.validate_dispatch();
+ return first;
+}
+
+TEST(RemoveTests, TestRemoveIfDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::remove_if(sys, vec.begin(), vec.end(), 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename ForwardIterator,
+ typename Predicate>
+__host__ __device__
+ForwardIterator remove_if(my_tag,
+ ForwardIterator first,
+ ForwardIterator,
+ Predicate)
+{
+ *first = 13;
+ return first;
+}
+
+TEST(RemoveTests, TestRemoveIfDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::remove_if(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ 0);
+
+ ASSERT_EQ(13, vec.front());
+}
+
+TYPED_TEST(RemoveTests, TestRemoveIfStencilSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector data(5);
+ data[0] = 1;
+ data[1] = 2;
+ data[2] = 1;
+ data[3] = 3;
+ data[4] = 2;
+
+ Vector stencil(5);
+ stencil[0] = 0;
+ stencil[1] = 1;
+ stencil[2] = 0;
+ stencil[3] = 0;
+ stencil[4] = 1;
+
+ typename Vector::iterator end = thrust::remove_if(data.begin(),
+ data.end(),
+ stencil.begin(),
+ thrust::identity<T>());
+
+ ASSERT_EQ(end - data.begin(), 3);
+
+ ASSERT_EQ(data[0], 1);
+ ASSERT_EQ(data[1], 1);
+ ASSERT_EQ(data[2], 3);
+}
+
+template<typename ForwardIterator,
+ typename InputIterator,
+ typename Predicate>
+__host__ __device__
+ForwardIterator remove_if(my_system &system,
+ ForwardIterator first,
+ ForwardIterator,
+ InputIterator,
+ Predicate)
+{
+ system.validate_dispatch();
+ return first;
+}
+
+TEST(RemoveTests, TestRemoveIfStencilDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::remove_if(sys,
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename ForwardIterator,
+ typename InputIterator,
+ typename Predicate>
+__host__ __device__
+ForwardIterator remove_if(my_tag,
+ ForwardIterator first,
+ ForwardIterator,
+ InputIterator,
+ Predicate)
+{
+ *first = 13;
+ return first;
+}
+
+TEST(RemoveTests, TestRemoveIfStencilDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::remove_if(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ 0);
+
+ ASSERT_EQ(13, vec.front());
+}
+
+TYPED_TEST(RemoveTests, TestRemoveCopyIfSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector data(5);
+ data[0] = 1;
+ data[1] = 2;
+ data[2] = 1;
+ data[3] = 3;
+ data[4] = 2;
+
+ Vector result(5);
+
+ typename Vector::iterator end = thrust::remove_copy_if(data.begin(),
+ data.end(),
+ result.begin(),
+ is_even<T>());
+
+ ASSERT_EQ(end - result.begin(), 3);
+
+ ASSERT_EQ(result[0], 1);
+ ASSERT_EQ(result[1], 1);
+ ASSERT_EQ(result[2], 3);
+}
+
+template<typename InputIterator,
+ typename OutputIterator,
+ typename Predicate>
+__host__ __device__
+InputIterator remove_copy_if(my_system &system,
+ InputIterator first,
+ InputIterator,
+ OutputIterator,
+ Predicate)
+{
+ system.validate_dispatch();
+ return first;
+}
+
+TEST(RemoveTests, TestRemoveCopyIfDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::remove_copy_if(sys,
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename InputIterator,
+ typename OutputIterator,
+ typename Predicate>
+__host__ __device__
+InputIterator remove_copy_if(my_tag,
+ InputIterator first,
+ InputIterator,
+ OutputIterator,
+ Predicate)
+{
+ *first = 13;
+ return first;
+}
+
+TEST(RemoveTests, TestRemoveCopyIfDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::remove_copy_if(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ 0);
+
+ ASSERT_EQ(13, vec.front());
+}
+
+TYPED_TEST(RemoveTests, TestRemoveCopyIfStencilSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector data(5);
+ data[0] = T(1);
+ data[1] = T(2);
+ data[2] = T(1);
+ data[3] = T(3);
+ data[4] = T(2);
+
+ Vector stencil(5);
+ stencil[0] = T(0);
+ stencil[1] = T(1);
+ stencil[2] = T(0);
+ stencil[3] = T(0);
+ stencil[4] = T(1);
+
+ Vector result(5);
+
+ typename Vector::iterator end = thrust::remove_copy_if(data.begin(),
+ data.end(),
+ stencil.begin(),
+ result.begin(),
+ thrust::identity<T>());
+
+ ASSERT_EQ(end - result.begin(), 3);
+
+ ASSERT_EQ(result[0], T(1));
+ ASSERT_EQ(result[1], T(1));
+ ASSERT_EQ(result[2], T(3));
+}
+
+template<typename InputIterator1,
+ typename InputIterator2,
+ typename OutputIterator,
+ typename Predicate>
+__host__ __device__
+OutputIterator remove_copy_if(my_system &system,
+ InputIterator1,
+ InputIterator1,
+ InputIterator2,
+ OutputIterator result,
+ Predicate)
+{
+ system.validate_dispatch();
+ return result;
+}
+
+TEST(RemoveTests, TestRemoveCopyIfStencilDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::remove_copy_if(sys,
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename InputIterator1,
+ typename InputIterator2,
+ typename OutputIterator,
+ typename Predicate>
+__host__ __device__
+OutputIterator remove_copy_if(my_tag,
+ InputIterator1,
+ InputIterator1,
+ InputIterator2,
+ OutputIterator result,
+ Predicate)
+{
+ *result = 13;
+ return result;
+}
+
+TEST(RemoveTests, TestRemoveCopyIfStencilDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::remove_copy_if(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ 0);
+
+ ASSERT_EQ(13, vec.front());
+}
+
+TYPED_TEST(RemoveTests, TestRemove)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_data = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ thrust::device_vector<T> d_data = h_data;
+
+ size_t h_size = thrust::remove(h_data.begin(), h_data.end(), T(0)) - h_data.begin();
+ size_t d_size = thrust::remove(d_data.begin(), d_data.end(), T(0)) - d_data.begin();
+
+ ASSERT_EQ(h_size, d_size);
+
+ h_data.resize(h_size);
+ d_data.resize(d_size);
+
+ ASSERT_EQ(h_data, d_data);
+ }
+}
+
+TYPED_TEST(RemoveTests, TestRemoveIf)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_data = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ thrust::device_vector<T> d_data = h_data;
+
+ size_t h_size = thrust::remove_if(h_data.begin(), h_data.end(), is_true<T>()) - h_data.begin();
+ size_t d_size = thrust::remove_if(d_data.begin(), d_data.end(), is_true<T>()) - d_data.begin();
+
+ ASSERT_EQ(h_size, d_size);
+
+ h_data.resize(h_size);
+ d_data.resize(d_size);
+
+ ASSERT_EQ(h_data, d_data);
+ }
+}
+
+TYPED_TEST(RemoveTests, TestRemoveIfStencil)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_data = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ thrust::device_vector<T> d_data = h_data;
+
+ thrust::host_vector<T> h_stencil = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+ thrust::device_vector<T> d_stencil = h_stencil;
+
+ size_t h_size = thrust::remove_if(h_data.begin(), h_data.end(), h_stencil.begin(), is_true<T>()) - h_data.begin();
+ size_t d_size = thrust::remove_if(d_data.begin(), d_data.end(), d_stencil.begin(), is_true<T>()) - d_data.begin();
+
+ ASSERT_EQ(h_size, d_size);
+
+ h_data.resize(h_size);
+ d_data.resize(d_size);
+
+ ASSERT_EQ(h_data, d_data);
+ }
+}
+
+TYPED_TEST(RemoveTests, TestRemoveCopy)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_data = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ thrust::device_vector<T> d_data = h_data;
+
+ thrust::host_vector<T> h_result(size);
+ thrust::device_vector<T> d_result(size);
+
+ size_t h_size = thrust::remove_copy(h_data.begin(), h_data.end(), h_result.begin(), T(0)) - h_result.begin();
+ size_t d_size = thrust::remove_copy(d_data.begin(), d_data.end(), d_result.begin(), T(0)) - d_result.begin();
+
+ ASSERT_EQ(h_size, d_size);
+
+ h_data.resize(h_size);
+ d_data.resize(d_size);
+
+ ASSERT_EQ(h_data, d_data);
+ }
+}
+
+TYPED_TEST(RemoveTests, TestRemoveCopyToDiscardIterator)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_data = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ thrust::device_vector<T> d_data = h_data;
+
+ size_t num_zeros = thrust::count(h_data.begin(), h_data.end(), T(0));
+ size_t num_nonzeros = h_data.size() - num_zeros;
+
+ thrust::discard_iterator<> h_result =
+ thrust::remove_copy(h_data.begin(), h_data.end(), thrust::make_discard_iterator(), T(0));
+
+ thrust::discard_iterator<> d_result =
+ thrust::remove_copy(d_data.begin(), d_data.end(), thrust::make_discard_iterator(), T(0));
+
+ thrust::discard_iterator<> reference(num_nonzeros);
+
+ ASSERT_EQ(reference, h_result);
+ ASSERT_EQ(reference, d_result);
+ }
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

ADD
test/test_transform_reduce.cpp
+// MIT License
+//
+// Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#include <iostream>
+#include <type_traits>
+#include <cstdlib>
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/transform_reduce.h>
+
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/iterator_traits.h>
+#include <thrust/iterator/retag.h>
+#include <thrust/device_vector.h>
+
+#include "test_utils.hpp"
+
+template<
+ class Input
+>
+struct Params
+{
+ using input_type = Input;
+};
+
+template<class Params>
+class TransformReduceTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<thrust::host_vector<short>>,
+ Params<thrust::host_vector<int>>,
+ Params<thrust::host_vector<long long>>,
+ Params<thrust::host_vector<unsigned short>>,
+ Params<thrust::host_vector<unsigned int>>,
+ Params<thrust::host_vector<unsigned long long>>,
+ Params<thrust::host_vector<float>>,
+ Params<thrust::host_vector<double>>,
+ Params<thrust::device_vector<short>>,
+ Params<thrust::device_vector<int>>,
+ Params<thrust::device_vector<long long>>,
+ Params<thrust::device_vector<unsigned short>>,
+ Params<thrust::device_vector<unsigned int>>,
+ Params<thrust::device_vector<unsigned long long>>,
+ Params<thrust::device_vector<float>>,
+ Params<thrust::device_vector<double>>
+> TransformReduceTestsParams;
+
+TYPED_TEST_CASE(TransformReduceTests, TransformReduceTestsParams);
+
+template<class Params>
+class TransformReduceIntegerTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<thrust::host_vector<short>>,
+ Params<thrust::host_vector<int>>,
+ Params<thrust::host_vector<long long>>,
+ Params<thrust::device_vector<short>>,
+ Params<thrust::device_vector<int>>,
+ Params<thrust::device_vector<long long>>
+> TransformReduceIntegerTestsParams;
+
+TYPED_TEST_CASE(TransformReduceIntegerTests, TransformReduceIntegerTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+/*template<typename InputIterator,
+ typename UnaryFunction,
+ typename OutputType,
+ typename BinaryFunction>
+OutputType transform_reduce(my_system &system,
+ InputIterator,
+ InputIterator,
+ UnaryFunction,
+ OutputType init,
+ BinaryFunction)
+{
+ system.validate_dispatch();
+ return init;
+}
+
+TEST(TransformReduceTests, TestTransformReduceDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::transform_reduce(sys,
+ vec.begin(),
+ vec.begin(),
+ 0,
+ 0,
+ 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename InputIterator,
+ typename UnaryFunction,
+ typename OutputType,
+ typename BinaryFunction>
+OutputType transform_reduce(my_tag,
+ InputIterator first,
+ InputIterator,
+ UnaryFunction,
+ OutputType init,
+ BinaryFunction)
+{
+ *first = 13;
+ return init;
+}
+
+TEST(TransformReduceTests, TestTransformReduceDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::transform_reduce(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ 0,
+ 0,
+ 0);
+
+ ASSERT_EQ(13, vec.front());
+}*/
+
+TYPED_TEST(TransformReduceTests, TestTransformReduceSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector data(3);
+ data[0] = T(1); data[1] = T(-2); data[2] = T(3);
+
+ T init = T(10);
+ T result = thrust::transform_reduce(data.begin(), data.end(), thrust::negate<T>(), init, thrust::plus<T>());
+
+ ASSERT_EQ(result, T(8));
+}
+
+TYPED_TEST(TransformReduceIntegerTests, TestTransformReduce)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_data = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ thrust::device_vector<T> d_data = h_data;
+
+ T init = T(13);
+
+ T cpu_result = thrust::transform_reduce(h_data.begin(), h_data.end(), thrust::negate<T>(), init, thrust::plus<T>());
+ T gpu_result = thrust::transform_reduce(d_data.begin(), d_data.end(), thrust::negate<T>(), init, thrust::plus<T>());
+
+ ASSERT_NEAR(cpu_result, gpu_result, std::abs(T(0.01 * cpu_result)) );
+ }
+}
+
+TYPED_TEST(TransformReduceIntegerTests, TestTransformReduceFromConst)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_data = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ thrust::device_vector<T> d_data = h_data;
+
+ T init = T(13);
+
+ T cpu_result = thrust::transform_reduce(h_data.cbegin(), h_data.cend(), thrust::negate<T>(), init, thrust::plus<T>());
+ T gpu_result = thrust::transform_reduce(d_data.cbegin(), d_data.cend(), thrust::negate<T>(), init, thrust::plus<T>());
+
+ ASSERT_NEAR(cpu_result, gpu_result, std::abs(T(0.01 * cpu_result)) );
+ }
+}
+
+TYPED_TEST(TransformReduceTests, TestTransformReduceCountingIterator)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+ using space = typename thrust::iterator_system<typename Vector::iterator>::type;
+ if( std::is_signed<T>::value )
+ {
+ thrust::counting_iterator<T, space> first(1);
+
+ T result = thrust::transform_reduce(first, first + 3, thrust::negate<short>(), 0, thrust::plus<short>());
+
+ ASSERT_EQ(result, -6);
+ }
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

MODIFY
test/test_vector_insert.cpp
test/test_vector_insert.cpp
@@ -222,33 +222,13 @@ TYPED_TEST(VectorInsertTests, TestVectorRangeInsert)
const std::vector<size_t> sizes = get_sizes();
for(auto size : sizes)
{
- thrust::host_vector<T> h_src;
- if (std::is_floating_point<T>::value)
- {
- h_src = get_random_data<T>(size + 3, (T)-1000, (T)+1000);
- }
- else
- {
- h_src = get_random_data<T>(
- size + 3,
+ thrust::host_vector<T> h_src = get_random_data<T>(size + 3,
std::numeric_limits<T>::min(),
- std::numeric_limits<T>::max()
- );
- }
+ std::numeric_limits<T>::max());
- thrust::host_vector<T> h_dst;
- if (std::is_floating_point<T>::value)
- {
- h_dst = get_random_data<T>(size, (T)-1000, (T)+1000);
- }
- else
- {
- h_dst = get_random_data<T>(
- size,
+ thrust::host_vector<T> h_dst = get_random_data<T>(size,
std::numeric_limits<T>::min(),
- std::numeric_limits<T>::max()
- );
- }
+ std::numeric_limits<T>::max());
thrust::device_vector<T> d_src = h_src;
thrust::device_vector<T> d_dst = h_dst;
@@ -403,19 +383,9 @@ TYPED_TEST(VectorInsertTests, TestVectorFillInsert)
const std::vector<size_t> sizes = get_sizes();
for(auto size : sizes)
{
- thrust::host_vector<T> h_dst;
- if (std::is_floating_point<T>::value)
- {
- h_dst = get_random_data<T>(size + 2, (T)-1000, (T)+1000);
- }
- else
- {
- h_dst = get_random_data<T>(
- size + 2,
+ thrust::host_vector<T> h_dst = get_random_data<T>(size + 2,
std::numeric_limits<T>::min(),
- std::numeric_limits<T>::max()
- );
- }
+ std::numeric_limits<T>::max());
thrust::device_vector<T> d_dst = h_dst;

MODIFY
thrust/system/detail/adl/transform_reduce.h
thrust/system/detail/adl/transform_reduce.h
#if 0
#include <thrust/system/cpp/detail/transform_reduce.h>
#include <thrust/system/cuda/detail/transform_reduce.h>
+#include <thrust/system/hip/detail/transform_reduce.h>
#include <thrust/system/omp/detail/transform_reduce.h>
#include <thrust/system/tbb/detail/transform_reduce.h>
#endif
#define __THRUST_DEVICE_SYSTEM_TRANSFORM_REDUCE_HEADER <__THRUST_DEVICE_SYSTEM_ROOT/detail/transform_reduce.h>
#include __THRUST_DEVICE_SYSTEM_TRANSFORM_REDUCE_HEADER
#undef __THRUST_DEVICE_SYSTEM_TRANSFORM_REDUCE_HEADER
-

MODIFY
thrust/system/hip/detail/transform_reduce.h
thrust/system/hip/detail/transform_reduce.h
-// STREAMHPC TODO implement
+/******************************************************************************
+ * Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ * * Neither the name of the NVIDIA CORPORATION nor the
+ * names of its contributors may be used to endorse or promote products
+ * derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+#pragma once
+
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <iterator>
+#include <thrust/system/hip/detail/reduce.h>
+#include <thrust/distance.h>
+
+BEGIN_NS_THRUST
+namespace hip_rocprim {
+
+template <class Derived,
+ class InputIt,
+ class TransformOp,
+ class T,
+ class ReduceOp>
+T __host__ __device__
+transform_reduce(execution_policy<Derived> &policy,
+ InputIt first,
+ InputIt last,
+ TransformOp transform_op,
+ T init,
+ ReduceOp reduce_op)
+{
+ typedef typename iterator_traits<InputIt>::difference_type size_type;
+ size_type num_items = static_cast<size_type>(thrust::distance(first, last));
+ typedef transform_input_iterator_t<T,
+ InputIt,
+ TransformOp>
+ transformed_iterator_t;
+
+ return hip_rocprim::reduce_n(policy,
+ transformed_iterator_t(first, transform_op),
+ num_items,
+ init,
+ reduce_op);
+}
+
+} // namespace hip_rocprim
+END_NS_THRUST
+#endif

MODIFY
thrust/system/hip/execution_policy.h
thrust/system/hip/execution_policy.h
#include <thrust/system/hip/detail/copy.h>
#include <thrust/system/hip/detail/copy_if.h>
// #include <thrust/system/hip/detail/count.h>
-// #include <thrust/system/hip/detail/equal.h>
+#include <thrust/system/hip/detail/equal.h>
// #include <thrust/system/hip/detail/extrema.h>
#include <thrust/system/hip/detail/fill.h>
-// #include <thrust/system/hip/detail/find.h>
+#include <thrust/system/hip/detail/find.h>
#include <thrust/system/hip/detail/for_each.h>
// #include <thrust/system/hip/detail/gather.h>
#include <thrust/system/hip/detail/generate.h>
// #include <thrust/system/hip/detail/inner_product.h>
-// #include <thrust/system/hip/detail/mismatch.h>
+#include <thrust/system/hip/detail/mismatch.h>
// #include <thrust/system/hip/detail/partition.h>
// #include <thrust/system/hip/detail/reduce_by_key.h>
#include <thrust/system/hip/detail/remove.h>
#include <thrust/system/hip/detail/replace.h>
// #include <thrust/system/hip/detail/reverse.h>
-// #include <thrust/system/hip/detail/scatter.h>
+#include <thrust/system/hip/detail/scatter.h>
#include <thrust/system/hip/detail/swap_ranges.h>
#include <thrust/system/hip/detail/tabulate.h>
#include <thrust/system/hip/detail/transform.h>
-// #include <thrust/system/hip/detail/transform_reduce.h>
+#include <thrust/system/hip/detail/transform_reduce.h>
// #include <thrust/system/hip/detail/transform_scan.h>
#include <thrust/system/hip/detail/uninitialized_copy.h>
#include <thrust/system/hip/detail/uninitialized_fill.h>
// // fail
// // ----------------
// // fails with mixed types
-// #include <thrust/system/hip/detail/reduce.h>
+#include <thrust/system/hip/detail/reduce.h>
// // mixed types are not compiling, commented in testing/scan.cu
// #include <thrust/system/hip/detail/scan.h>

Implement transform reduce and remove tests

Updated the execution policy header too.
Replaced the random sample in vector insert tests.