Add uninitialized_fill, and counting_iterator, tabulate, distance tests
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -58,15 +58,19 @@ endfunction()
add_hipthrust_test("thrust.hip.allocator" test_allocator.cpp)
add_hipthrust_test("thrust.hip.copy" test_copy.cpp)
+add_hipthrust_test("thrust.hip.counting_iterator" test_counting_iterator.cpp)
add_hipthrust_test("thrust.hip.device_delete" test_device_delete.cpp)
add_hipthrust_test("thrust.hip.device_ptr" test_device_ptr.cpp)
add_hipthrust_test("thrust.hip.device_reference" test_device_reference.cpp)
+add_hipthrust_test("thrust.hip.distance" test_distance.cpp)
add_hipthrust_test("thrust.hip.fill" test_fill.cpp)
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.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.uninitialized_fill" test_uninitialized_fill.cpp)
add_hipthrust_test("thrust.hip.vector" test_vector.cpp)
add_hipthrust_test("thrust.hip.vector_insert" test_vector_insert.cpp)
add_hipthrust_test("thrust.hip.vector_manipulation" test_vector_manipulation.cpp)

ADD
test/test_counting_iterator.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 <vector>
+#include <list>
+#include <limits>
+#include <utility>
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/sort.h>
+// TODO: implement binary search
+//#include <thrust/binary_search.h>
+#include <thrust/distance.h>
+#include <thrust/detail/cstdint.h>
+
+// HIP API
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <hip/hip_runtime_api.h>
+#include <hip/hip_runtime.h>
+
+#define HIP_CHECK(condition) ASSERT_EQ(condition, hipSuccess)
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+TEST(CountingIteratorTests, TestCountingIteratorCopyConstructor)
+{
+ thrust::counting_iterator<int> iter0(100);
+
+ thrust::counting_iterator<int> iter1(iter0);
+
+ ASSERT_EQ(iter0, iter1);
+ ASSERT_EQ(*iter0, *iter1);
+
+ // construct from related space
+ thrust::counting_iterator<int, thrust::host_system_tag> h_iter = iter0;
+ ASSERT_EQ(*iter0, *h_iter);
+
+ thrust::counting_iterator<int, thrust::device_system_tag> d_iter = iter0;
+ ASSERT_EQ(*iter0, *d_iter);
+}
+
+TEST(CountingIteratorTests, TestCountingIteratorIncrement)
+{
+ thrust::counting_iterator<int> iter(0);
+
+ ASSERT_EQ(*iter, 0);
+
+ iter++;
+
+ ASSERT_EQ(*iter, 1);
+
+ iter++;
+ iter++;
+
+ ASSERT_EQ(*iter, 3);
+
+ iter += 5;
+
+ ASSERT_EQ(*iter, 8);
+
+ iter -= 10;
+
+ ASSERT_EQ(*iter, -2);
+}
+
+TEST(CountingIteratorTests, TestCountingIteratorComparison)
+{
+ thrust::counting_iterator<int> iter1(0);
+ thrust::counting_iterator<int> iter2(0);
+
+ ASSERT_EQ(iter1 - iter2, 0);
+ ASSERT_EQ(iter1 == iter2, true);
+
+ iter1++;
+
+ ASSERT_EQ(iter1 - iter2, 1);
+ ASSERT_EQ(iter1 == iter2, false);
+
+ iter2++;
+
+ ASSERT_EQ(iter1 - iter2, 0);
+ ASSERT_EQ(iter1 == iter2, true);
+
+ iter1 += 100;
+ iter2 += 100;
+
+ ASSERT_EQ(iter1 - iter2, 0);
+ ASSERT_EQ(iter1 == iter2, true);
+}
+
+TEST(CountingIteratorTests, TestCountingIteratorFloatComparison)
+{
+ thrust::counting_iterator<float> iter1(0);
+ thrust::counting_iterator<float> iter2(0);
+
+ ASSERT_EQ(iter1 - iter2, 0);
+ ASSERT_EQ(iter1 == iter2, true);
+ ASSERT_EQ(iter1 < iter2, false);
+ ASSERT_EQ(iter2 < iter1, false);
+
+ iter1++;
+
+ ASSERT_EQ(iter1 - iter2, 1);
+ ASSERT_EQ(iter1 == iter2, false);
+ ASSERT_EQ(iter2 < iter1, true);
+ ASSERT_EQ(iter1 < iter2, false);
+
+ iter2++;
+
+ ASSERT_EQ(iter1 - iter2, 0);
+ ASSERT_EQ(iter1 == iter2, true);
+ ASSERT_EQ(iter1 < iter2, false);
+ ASSERT_EQ(iter2 < iter1, false);
+
+ iter1 += 100;
+ iter2 += 100;
+
+ ASSERT_EQ(iter1 - iter2, 0);
+ ASSERT_EQ(iter1 == iter2, true);
+ ASSERT_EQ(iter1 < iter2, false);
+ ASSERT_EQ(iter2 < iter1, false);
+
+
+ thrust::counting_iterator<float> iter3(0);
+ thrust::counting_iterator<float> iter4(0.5);
+
+ ASSERT_EQ(iter3 - iter4, 0);
+ ASSERT_EQ(iter3 == iter4, true);
+ ASSERT_EQ(iter3 < iter4, false);
+ ASSERT_EQ(iter4 < iter3, false);
+
+ iter3++; // iter3 = 1.0, iter4 = 0.5
+
+ ASSERT_EQ(iter3 - iter4, 0);
+ ASSERT_EQ(iter3 == iter4, true);
+ ASSERT_EQ(iter3 < iter4, false);
+ ASSERT_EQ(iter4 < iter3, false);
+
+ iter4++; // iter3 = 1.0, iter4 = 1.5
+
+ ASSERT_EQ(iter3 - iter4, 0);
+ ASSERT_EQ(iter3 == iter4, true);
+ ASSERT_EQ(iter3 < iter4, false);
+ ASSERT_EQ(iter4 < iter3, false);
+
+ iter4++; // iter3 = 1.0, iter4 = 2.5
+
+ ASSERT_EQ(iter3 - iter4, -1);
+ ASSERT_EQ(iter4 - iter3, 1);
+ ASSERT_EQ(iter3 == iter4, false);
+ ASSERT_EQ(iter3 < iter4, true);
+ ASSERT_EQ(iter4 < iter3, false);
+}
+
+TEST(CountingIteratorTests, TestCountingIteratorDistance)
+{
+ thrust::counting_iterator<int> iter1(0);
+ thrust::counting_iterator<int> iter2(5);
+
+ ASSERT_EQ(thrust::distance(iter1, iter2), 5);
+
+ iter1++;
+
+ ASSERT_EQ(thrust::distance(iter1, iter2), 4);
+
+ iter2 += 100;
+
+ ASSERT_EQ(thrust::distance(iter1, iter2), 104);
+
+ iter2 += 1000;
+
+ ASSERT_EQ(thrust::distance(iter1, iter2), 1104);
+}
+
+TEST(CountingIteratorTests, TestCountingIteratorUnsignedType)
+{
+ thrust::counting_iterator<unsigned int> iter0(0);
+ thrust::counting_iterator<unsigned int> iter1(5);
+
+ ASSERT_EQ(iter1 - iter0, 5);
+ ASSERT_EQ(iter0 - iter1, -5);
+ ASSERT_EQ(iter0 != iter1, true);
+ ASSERT_EQ(iter0 < iter1, true);
+ ASSERT_EQ(iter1 < iter0, false);
+}
+
+// TODO: implement binary search
+/*TEST(CountingIteratorTests, TestCountingIteratorLowerBound)
+{
+ size_t n = 10000;
+ const size_t M = 100;
+
+ thrust::host_vector<unsigned int> h_data = unittest::random_integers<unsigned int>(n);
+ for(unsigned int i = 0; i < n; ++i)
+ h_data[i] %= M;
+
+ thrust::sort(h_data.begin(), h_data.end());
+
+ thrust::device_vector<unsigned int> d_data = h_data;
+
+ thrust::counting_iterator<unsigned int> search_begin(0);
+ thrust::counting_iterator<unsigned int> search_end(M);
+
+
+ thrust::host_vector<unsigned int> h_result(M);
+ thrust::device_vector<unsigned int> d_result(M);
+
+
+ thrust::lower_bound(h_data.begin(), h_data.end(), search_begin, search_end, h_result.begin());
+
+ thrust::lower_bound(d_data.begin(), d_data.end(), search_begin, search_end, d_result.begin());
+
+ ASSERT_EQ(h_result, d_result);
+}*/
+
+TEST(CountingIteratorTests, TestCountingIteratorDifference)
+{
+ using Iterator = typename thrust::counting_iterator<thrust::detail::uint64_t>;
+ using Difference = typename thrust::iterator_difference<Iterator>::type;
+
+ Difference diff = std::numeric_limits<thrust::detail::uint32_t>::max() + 1;
+
+ Iterator first(0);
+ Iterator last = first + diff;
+
+ ASSERT_EQ(diff, last - first);
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

ADD
test/test_distance.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 <vector>
+#include <list>
+#include <limits>
+#include <utility>
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/distance.h>
+#include <thrust/host_vector.h>
+#include <thrust/device_vector.h>
+
+// HIP API
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <hip/hip_runtime_api.h>
+#include <hip/hip_runtime.h>
+
+#define HIP_CHECK(condition) ASSERT_EQ(condition, hipSuccess)
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class DistanceTests : 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>>
+> DistanceTestsParams;
+
+TYPED_TEST_CASE(DistanceTests, DistanceTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+
+TYPED_TEST(DistanceTests, TestDistance)
+{
+ using Vector = typename TestFixture::input_type;
+ using Iterator = typename Vector::iterator;
+
+ Vector v(100);
+
+ Iterator i = v.begin();
+
+ ASSERT_EQ(thrust::distance(i, v.end()), 100);
+
+ i++;
+
+ ASSERT_EQ(thrust::distance(i, v.end()), 99);
+
+ i += 49;
+
+ ASSERT_EQ(thrust::distance(i, v.end()), 50);
+
+ ASSERT_EQ(thrust::distance(i, i), 0);
+}
+
+TYPED_TEST(DistanceTests, TestDistanceLarge)
+{
+ using Vector = typename TestFixture::input_type;
+ using Iterator = typename Vector::iterator;
+
+ Vector v(1000);
+
+ Iterator i = v.begin();
+
+ ASSERT_EQ(thrust::distance(i, v.end()), 1000);
+
+ i++;
+
+ ASSERT_EQ(thrust::distance(i, v.end()), 999);
+
+ i += 49;
+
+ ASSERT_EQ(thrust::distance(i, v.end()), 950);
+
+ i += 950;
+
+ ASSERT_EQ(thrust::distance(i, v.end()), 0);
+
+ ASSERT_EQ(thrust::distance(i, i), 0);
+}
+
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

ADD
test/test_tabulate.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 <vector>
+#include <list>
+#include <limits>
+#include <utility>
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/tabulate.h>
+#include <thrust/functional.h>
+#include <thrust/iterator/discard_iterator.h>
+#include <thrust/iterator/retag.h>
+#include <thrust/device_vector.h>
+
+// HIP API
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <hip/hip_runtime_api.h>
+#include <hip/hip_runtime.h>
+
+#define HIP_CHECK(condition) ASSERT_EQ(condition, hipSuccess)
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+#include "test_utils.hpp"
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class TabulateTests : 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>>
+> TabulateTestsParams;
+
+TYPED_TEST_CASE(TabulateTests, TabulateTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+template<typename ForwardIterator, typename UnaryOperation>
+void tabulate(my_system &system, ForwardIterator, ForwardIterator, UnaryOperation)
+{
+ system.validate_dispatch();
+}
+
+TEST(TabulateTests, TestTabulateDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::tabulate(sys, vec.begin(), vec.end(), thrust::identity<int>());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename ForwardIterator, typename UnaryOperation>
+void tabulate(my_tag, ForwardIterator first, ForwardIterator, UnaryOperation)
+{
+ *first = 13;
+}
+
+TEST(TabulateTests, TestTabulateDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::tabulate(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.end()),
+ thrust::identity<int>());
+
+ ASSERT_EQ(13, vec.front());
+}
+
+TYPED_TEST(TabulateTests, TestTabulateSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+ using namespace thrust::placeholders;
+
+ Vector v(5);
+
+ thrust::tabulate(v.begin(), v.end(), thrust::identity<T>());
+
+ ASSERT_EQ(v[0], T(0));
+ ASSERT_EQ(v[1], T(1));
+ ASSERT_EQ(v[2], T(2));
+ ASSERT_EQ(v[3], T(3));
+ ASSERT_EQ(v[4], T(4));
+
+ thrust::tabulate(v.begin(), v.end(), -_1);
+
+ ASSERT_EQ(v[0], T( 0));
+ ASSERT_EQ(v[1], T(-1));
+ ASSERT_EQ(v[2], T(-2));
+ ASSERT_EQ(v[3], T(-3));
+ ASSERT_EQ(v[4], T(-4));
+
+ thrust::tabulate(v.begin(), v.end(), _1 * _1 * _1);
+
+ ASSERT_EQ(v[0], T(0));
+ ASSERT_EQ(v[1], T(1));
+ ASSERT_EQ(v[2], T(8));
+ ASSERT_EQ(v[3], T(27));
+ ASSERT_EQ(v[4], T(64));
+}
+
+TYPED_TEST(TabulateTests, TestTabulate)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+ using namespace thrust::placeholders;
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ thrust::host_vector<T> h_data(size);
+ thrust::device_vector<T> d_data(size);
+
+ thrust::tabulate(h_data.begin(), h_data.end(), _1 * _1 + T(13));
+ thrust::tabulate(d_data.begin(), d_data.end(), _1 * _1 + T(13));
+
+ for (size_t i = 0; i < size; i++)
+ {
+ ASSERT_EQ(h_data[i], d_data[i]) << "where index = " << i;
+ }
+
+ thrust::tabulate(h_data.begin(), h_data.end(), (_1 - T(7)) * _1);
+ thrust::tabulate(d_data.begin(), d_data.end(), (_1 - T(7)) * _1);
+
+ ASSERT_EQ(h_data, d_data);
+ }
+}
+
+TEST(TabulateTests, TestTabulateToDiscardIterator)
+{
+ for (auto size : get_sizes())
+ {
+ thrust::tabulate(thrust::discard_iterator<thrust::device_system_tag>(),
+ thrust::discard_iterator<thrust::device_system_tag>(size),
+ thrust::identity<int>());
+ }
+ // nothing to check -- just make sure it compiles
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

ADD
test/test_uninitialized_fill.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.
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/uninitialized_fill.h>
+#include <thrust/device_malloc_allocator.h>
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+#include <thrust/iterator/retag.h>
+
+// HIP API
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <hip/hip_runtime_api.h>
+#include <hip/hip_runtime.h>
+
+#define HIP_CHECK(condition) ASSERT_EQ(condition, hipSuccess)
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+#include "test_utils.hpp"
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class UninitializedFillTests : 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>>
+> UninitializedFillTestsParams;
+
+TYPED_TEST_CASE(UninitializedFillTests, UninitializedFillTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+template<typename ForwardIterator, typename T>
+void uninitialized_fill(my_system &system,
+ ForwardIterator,
+ ForwardIterator,
+ const T &)
+{
+ system.validate_dispatch();
+}
+
+TEST(UninitializedFillTests, TestUninitializedFillDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::uninitialized_fill(sys, vec.begin(), vec.begin(), 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename ForwardIterator, typename T>
+void uninitialized_fill(my_tag,
+ ForwardIterator first,
+ ForwardIterator,
+ const T &)
+{
+ *first = 13;
+}
+
+TEST(UninitializedFillTests, TestUninitializedFillDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::uninitialized_fill(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ 0);
+
+ ASSERT_EQ(13, vec.front());
+}
+
+template<typename ForwardIterator, typename Size, typename T>
+ForwardIterator uninitialized_fill_n(my_system &system,
+ ForwardIterator first,
+ Size,
+ const T &)
+{
+ system.validate_dispatch();
+ return first;
+}
+
+TEST(UninitializedFillTests, TestUninitializedFillNDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::uninitialized_fill_n(sys, vec.begin(), vec.size(), 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename ForwardIterator, typename Size, typename T>
+ForwardIterator uninitialized_fill_n(my_tag,
+ ForwardIterator first,
+ Size,
+ const T &)
+{
+ *first = 13;
+ return first;
+}
+
+TEST(UninitializedFillTests, TestUninitializedFillNDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::uninitialized_fill_n(sys,
+ vec.begin(),
+ vec.size(),
+ 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+TYPED_TEST(UninitializedFillTests, TestUninitializedFillPOD)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector v(5);
+ v[0] = T(0); v[1] = T(1); v[2] = T(2); v[3] = T(3); v[4] = T(4);
+
+ T exemplar(7);
+
+ thrust::uninitialized_fill(v.begin() + 1, v.begin() + 4, exemplar);
+
+ ASSERT_EQ(v[0], T(0));
+ ASSERT_EQ(v[1], exemplar);
+ ASSERT_EQ(v[2], exemplar);
+ ASSERT_EQ(v[3], exemplar);
+ ASSERT_EQ(v[4], T(4));
+
+ exemplar = T(8);
+
+ thrust::uninitialized_fill(v.begin() + 0, v.begin() + 3, exemplar);
+
+ ASSERT_EQ(v[0], exemplar);
+ ASSERT_EQ(v[1], exemplar);
+ ASSERT_EQ(v[2], exemplar);
+ ASSERT_EQ(v[3], T(7));
+ ASSERT_EQ(v[4], T(4));
+
+ exemplar = T(9);
+
+ thrust::uninitialized_fill(v.begin() + 2, v.end(), exemplar);
+
+ ASSERT_EQ(v[0], T(8));
+ ASSERT_EQ(v[1], T(8));
+ ASSERT_EQ(v[2], exemplar);
+ ASSERT_EQ(v[3], exemplar);
+ ASSERT_EQ(v[4], T(9));
+
+ exemplar = T(1);
+
+ thrust::uninitialized_fill(v.begin(), v.end(), exemplar);
+
+ ASSERT_EQ(v[0], exemplar);
+ ASSERT_EQ(v[1], exemplar);
+ ASSERT_EQ(v[2], exemplar);
+ ASSERT_EQ(v[3], exemplar);
+ ASSERT_EQ(v[4], exemplar);
+}
+
+struct CopyConstructTest
+{
+ CopyConstructTest(void)
+ :copy_constructed_on_host(false),
+ copy_constructed_on_device(false)
+ {}
+
+ __host__ __device__
+ CopyConstructTest(const CopyConstructTest&)
+ {
+#if defined(THRUST_HIP_DEVICE_CODE)
+ copy_constructed_on_device = true;
+ copy_constructed_on_host = false;
+#else
+ copy_constructed_on_device = false;
+ copy_constructed_on_host = true;
+#endif
+ }
+
+ __host__ __device__
+ CopyConstructTest &operator=(const CopyConstructTest &x)
+ {
+ copy_constructed_on_host = x.copy_constructed_on_host;
+ copy_constructed_on_device = x.copy_constructed_on_device;
+ return *this;
+ }
+
+ bool copy_constructed_on_host;
+ bool copy_constructed_on_device;
+};
+
+TEST(UninitializedFillTests, TestUninitializedFillNonPOD)
+{
+ using T = CopyConstructTest;
+ thrust::device_ptr<T> v = thrust::device_malloc<T>(5);
+
+ T exemplar;
+ ASSERT_EQ(false, exemplar.copy_constructed_on_device);
+ ASSERT_EQ(false, exemplar.copy_constructed_on_host);
+
+ T host_copy_of_exemplar(exemplar);
+ ASSERT_EQ(false, host_copy_of_exemplar.copy_constructed_on_device);
+ ASSERT_EQ(true, host_copy_of_exemplar.copy_constructed_on_host);
+
+ // copy construct v from the exemplar
+ thrust::uninitialized_fill(v, v + 1, exemplar);
+
+ T x;
+ ASSERT_EQ(false, x.copy_constructed_on_device);
+ ASSERT_EQ(false, x.copy_constructed_on_host);
+
+ x = v[0];
+ ASSERT_EQ(true, x.copy_constructed_on_device);
+ ASSERT_EQ(false, x.copy_constructed_on_host);
+
+ thrust::device_free(v);
+}
+
+TYPED_TEST(UninitializedFillTests, TestUninitializedFillNPOD)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector v(5);
+ v[0] = T(0); v[1] = T(1); v[2] = T(2); v[3] = T(3); v[4] = T(4);
+
+ T exemplar(7);
+
+ using Iterator = typename Vector::iterator;
+ Iterator iter = thrust::uninitialized_fill_n(v.begin() + 1, size_t(3), exemplar);
+
+ ASSERT_EQ(v[0], T(0));
+ ASSERT_EQ(v[1], exemplar);
+ ASSERT_EQ(v[2], exemplar);
+ ASSERT_EQ(v[3], exemplar);
+ ASSERT_EQ(v[4], T(4));
+ ASSERT_EQ(v.begin() + 4, iter);
+
+ exemplar = T(8);
+
+ iter = thrust::uninitialized_fill_n(v.begin() + 0, size_t(3), exemplar);
+
+ ASSERT_EQ(v[0], exemplar);
+ ASSERT_EQ(v[1], exemplar);
+ ASSERT_EQ(v[2], exemplar);
+ ASSERT_EQ(v[3], T(7));
+ ASSERT_EQ(v[4], T(4));
+ ASSERT_EQ(v.begin() + 3, iter);
+
+ exemplar = T(9);
+
+ iter = thrust::uninitialized_fill_n(v.begin() + 2, size_t(3), exemplar);
+
+ ASSERT_EQ(v[0], T(8));
+ ASSERT_EQ(v[1], T(8));
+ ASSERT_EQ(v[2], exemplar);
+ ASSERT_EQ(v[3], exemplar);
+ ASSERT_EQ(v[4], T(9));
+ ASSERT_EQ(v.end(), iter);
+
+ exemplar = T(1);
+
+ iter = thrust::uninitialized_fill_n(v.begin(), v.size(), exemplar);
+
+ ASSERT_EQ(v[0], exemplar);
+ ASSERT_EQ(v[1], exemplar);
+ ASSERT_EQ(v[2], exemplar);
+ ASSERT_EQ(v[3], exemplar);
+ ASSERT_EQ(v[4], exemplar);
+ ASSERT_EQ(v.end(), iter);
+}
+
+TEST(UninitializedFillTests, TestUninitializedFillNNonPOD)
+{
+ using T = CopyConstructTest;
+ thrust::device_ptr<T> v = thrust::device_malloc<T>(5);
+
+ T exemplar;
+ ASSERT_EQ(false, exemplar.copy_constructed_on_device);
+ ASSERT_EQ(false, exemplar.copy_constructed_on_host);
+
+ T host_copy_of_exemplar(exemplar);
+ ASSERT_EQ(false, host_copy_of_exemplar.copy_constructed_on_device);
+ ASSERT_EQ(true, host_copy_of_exemplar.copy_constructed_on_host);
+
+ // copy construct v from the exemplar
+ thrust::uninitialized_fill_n(v, size_t(1), exemplar);
+
+ T x;
+ ASSERT_EQ(false, x.copy_constructed_on_device);
+ ASSERT_EQ(false, x.copy_constructed_on_host);
+
+ x = v[0];
+ ASSERT_EQ(true, x.copy_constructed_on_device);
+ ASSERT_EQ(false, x.copy_constructed_on_host);
+
+ thrust::device_free(v);
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

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

MODIFY
thrust/system/hip/detail/uninitialized_fill.h
thrust/system/hip/detail/uninitialized_fill.h
*
******************************************************************************/
#pragma once
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <iterator>
+#include <thrust/distance.h>
+#include <thrust/system/hip/detail/execution_policy.h>
+#include <thrust/system/hip/detail/util.h>
+#include <thrust/system/hip/detail/parallel_for.h>
+
+BEGIN_NS_THRUST
+
+namespace hip_rocprim {
+
+namespace __uninitialized_fill {
+
+ template <class Iterator, class T>
+ struct functor
+ {
+ Iterator items;
+ T value;
+
+ typedef typename iterator_traits<Iterator>::value_type value_type;
+
+ THRUST_HIP_FUNCTION
+ functor(Iterator items_, T const& value_)
+ : items(items_), value(value_) {}
+
+ template<class Size>
+ void THRUST_HIP_DEVICE_FUNCTION operator()(Size idx)
+ {
+ value_type& out = raw_reference_cast(items[idx]);
+
+ ::new (static_cast<void *>(&out)) value_type(value);
+ }
+ }; // struct functor
+
+} // namespace __uninitialized_copy
+
+template <class Derived,
+ class Iterator,
+ class Size,
+ class T>
+Iterator __host__ __device__
+uninitialized_fill_n(execution_policy<Derived>& policy,
+ Iterator first,
+ Size count,
+ T const& x)
+{
+ typedef __uninitialized_fill::functor<Iterator,T> functor_t;
+
+ hip_rocprim::parallel_for(policy,
+ functor_t(first, x),
+ count);
+ return first + count;
+}
+
+template <class Derived,
+ class Iterator,
+ class T>
+void __host__ __device__
+uninitialized_fill(execution_policy<Derived>& policy,
+ Iterator first,
+ Iterator last,
+ T const& x)
+{
+ hip_rocprim::uninitialized_fill_n(policy,
+ first,
+ thrust::distance(first, last),
+ x);
+}
+
+} // namespace hip_rocprim
+
+END_NS_THRUST
+#endif

Add uninitialized_fill, and counting_iterator, tabulate, distance tests