Add vector insert and scatter tests
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -61,7 +61,9 @@ 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.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.transform" test_transform.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_scatter.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>
+#include <algorithm>
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/scatter.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/discard_iterator.h>
+#include <thrust/iterator/retag.h>
+#include <thrust/sequence.h>
+#include <thrust/device_vector.h>
+#include <thrust/fill.h>
+
+// HIP API
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <hip/hip_runtime_api.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 ScatterTests : 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>>
+> ScatterTestsParams;
+
+TYPED_TEST_CASE(ScatterTests, ScatterTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+TYPED_TEST(ScatterTests, TestScatterSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector map(5); // scatter indices
+ Vector src(5); // source vector
+ Vector dst(8); // destination vector
+
+ map[0] = T(6); map[1] = T(3); map[2] = T(1); map[3] = T(7); map[4] = T(2);
+ src[0] = T(0); src[1] = T(1); src[2] = T(2); src[3] = T(3); src[4] = T(4);
+ dst[0] = T(0); dst[1] = T(0); dst[2] = T(0); dst[3] = T(0); dst[4] = T(0); dst[5] = T(0); dst[6] = T(0); dst[7] = T(0);
+
+ thrust::scatter(src.begin(), src.end(), map.begin(), dst.begin());
+
+ ASSERT_EQ(dst[0], T(0));
+ ASSERT_EQ(dst[1], T(2));
+ ASSERT_EQ(dst[2], T(4));
+ ASSERT_EQ(dst[3], T(1));
+ ASSERT_EQ(dst[4], T(0));
+ ASSERT_EQ(dst[5], T(0));
+ ASSERT_EQ(dst[6], T(0));
+ ASSERT_EQ(dst[7], T(3));
+}
+
+template<typename InputIterator1,
+ typename InputIterator2,
+ typename RandomAccessIterator>
+void scatter(my_system &system,
+ InputIterator1,
+ InputIterator1,
+ InputIterator2,
+ RandomAccessIterator)
+{
+ system.validate_dispatch();
+}
+
+TEST(ScatterTests, TestScatterDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::scatter(sys,
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ vec.begin());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename InputIterator1,
+ typename InputIterator2,
+ typename RandomAccessIterator>
+void scatter(my_tag,
+ InputIterator1,
+ InputIterator1,
+ InputIterator2,
+ RandomAccessIterator output)
+{
+ *output = 13;
+}
+
+TEST(ScatterTests, TestScatterDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::scatter(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()));
+
+ ASSERT_EQ(13, vec.front());
+}
+
+TYPED_TEST(ScatterTests, TestScatter)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ const size_t output_size = std::min((size_t) 10, 2 * size);
+
+ thrust::host_vector<T> h_input(size, (T) 1);
+ thrust::device_vector<T> d_input(size, (T) 1);
+
+ thrust::host_vector<unsigned int> h_map= get_random_data<unsigned int>(size,
+ std::numeric_limits<unsigned int>::min(),
+ std::numeric_limits<unsigned int>::max());
+
+ for(size_t i = 0; i < size; i++)
+ h_map[i] = h_map[i] % output_size;
+
+ thrust::device_vector<unsigned int> d_map = h_map;
+
+ thrust::host_vector<T> h_output(output_size, T(0));
+ thrust::device_vector<T> d_output(output_size, T(0));
+
+ thrust::scatter(h_input.begin(), h_input.end(), h_map.begin(), h_output.begin());
+ thrust::scatter(d_input.begin(), d_input.end(), d_map.begin(), d_output.begin());
+
+ ASSERT_EQ(h_output, d_output);
+ }
+}
+
+TYPED_TEST(ScatterTests, TestScatterToDiscardIterator)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ const size_t output_size = std::min((size_t) 10, 2 * size);
+
+ thrust::host_vector<T> h_input(size, (T) 1);
+ thrust::device_vector<T> d_input(size, (T) 1);
+
+ thrust::host_vector<unsigned int> h_map= get_random_data<unsigned int>(size,
+ std::numeric_limits<unsigned int>::min(),
+ std::numeric_limits<unsigned int>::max());
+
+ for(size_t i = 0; i < size; i++)
+ h_map[i] = h_map[i] % output_size;
+
+ thrust::device_vector<unsigned int> d_map = h_map;
+
+ thrust::scatter(h_input.begin(), h_input.end(), h_map.begin(), thrust::make_discard_iterator());
+ thrust::scatter(d_input.begin(), d_input.end(), d_map.begin(), thrust::make_discard_iterator());
+
+ // there's nothing to check -- just make sure it compiles
+ }
+}
+
+TYPED_TEST(ScatterTests, TestScatterIfSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector flg(5); // predicate array
+ Vector map(5); // scatter indices
+ Vector src(5); // source vector
+ Vector dst(8); // destination vector
+
+ flg[0] = T(0); flg[1] = T(1); flg[2] = T(0); flg[3] = T(1); flg[4] = T(0);
+ map[0] = T(6); map[1] = T(3); map[2] = T(1); map[3] = T(7); map[4] = T(2);
+ src[0] = T(0); src[1] = T(1); src[2] = T(2); src[3] = T(3); src[4] = T(4);
+ dst[0] = T(0); dst[1] = T(0); dst[2] = T(0); dst[3] = T(0); dst[4] = T(0); dst[5] = T(0); dst[6] = T(0); dst[7] = T(0);
+
+ thrust::scatter_if(src.begin(), src.end(), map.begin(), flg.begin(), dst.begin());
+
+ ASSERT_EQ(dst[0], T(0));
+ ASSERT_EQ(dst[1], T(0));
+ ASSERT_EQ(dst[2], T(0));
+ ASSERT_EQ(dst[3], T(1));
+ ASSERT_EQ(dst[4], T(0));
+ ASSERT_EQ(dst[5], T(0));
+ ASSERT_EQ(dst[6], T(0));
+ ASSERT_EQ(dst[7], T(3));
+}
+
+template<typename InputIterator1,
+ typename InputIterator2,
+ typename InputIterator3,
+ typename RandomAccessIterator>
+void scatter_if(my_system &system,
+ InputIterator1,
+ InputIterator1,
+ InputIterator2,
+ InputIterator3,
+ RandomAccessIterator)
+{
+ system.validate_dispatch();
+}
+
+TEST(ScatterTests, TestScatterIfDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::scatter_if(sys,
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ vec.begin());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

MODIFY
test/test_utils.hpp
test/test_utils.hpp
#include <vector>
#include <thrust/host_vector.h>
+#include <thrust/execution_policy.h>
std::vector<size_t> get_sizes()
{
@@ -70,3 +71,40 @@ struct custom_compare_less
return lhs < rhs;
}
}; // end less
+
+class my_system : public thrust::device_execution_policy<my_system>
+{
+ public:
+ my_system(int)
+ : correctly_dispatched(false),
+ num_copies(0)
+ {}
+
+ my_system(const my_system &other)
+ : correctly_dispatched(false),
+ num_copies(other.num_copies + 1)
+ {}
+
+ void validate_dispatch()
+ {
+ correctly_dispatched = (num_copies == 0);
+ }
+
+ bool is_valid()
+ {
+ return correctly_dispatched;
+ }
+
+ private:
+ bool correctly_dispatched;
+
+ // count the number of copies so that we can validate
+ // that dispatch does not introduce any
+ unsigned int num_copies;
+
+
+ // disallow default construction
+ my_system();
+};
+
+struct my_tag : thrust::device_execution_policy<my_tag> {};

ADD
test/test_vector_insert.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/memory.h>
+#include <thrust/transform.h>
+// STREAMHPC TODO replace <thrust/detail/seq.h> with <thrust/execution_policy.h>
+#include <thrust/sequence.h>
+#include <thrust/device_malloc_allocator.h>
+#include <thrust/device_vector.h>
+#include <thrust/host_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 VectorInsertTests : 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>>
+> VectorInsertTestsParams;
+
+TYPED_TEST_CASE(VectorInsertTests, VectorInsertTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+TYPED_TEST(VectorInsertTests, TestVectorRangeInsertSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector v1(5);
+ thrust::sequence(v1.begin(), v1.end());
+
+ // test when insertion range fits inside capacity
+ // and the size of the insertion is greater than the number
+ // of displaced elements
+ Vector v2(3);
+ v2.reserve(10);
+ thrust::sequence(v2.begin(), v2.end());
+
+ size_t new_size = v2.size() + v1.size();
+ size_t insertion_size = v1.end() - v1.begin();
+ size_t num_displaced = v2.end() - (v2.begin() + 1);
+
+ ASSERT_EQ(true, v2.capacity() >= new_size);
+ ASSERT_EQ(true, insertion_size > num_displaced);
+
+ v2.insert(v2.begin() + 1,
+ v1.begin(), v1.end());
+
+ ASSERT_EQ(T(0), v2[0]);
+
+ ASSERT_EQ(T(0), v2[1]);
+ ASSERT_EQ(T(1), v2[2]);
+ ASSERT_EQ(T(2), v2[3]);
+ ASSERT_EQ(T(3), v2[4]);
+ ASSERT_EQ(T(4), v2[5]);
+
+ ASSERT_EQ(T(1), v2[6]);
+ ASSERT_EQ(T(2), v2[7]);
+
+ ASSERT_EQ(T(8), v2.size());
+ ASSERT_EQ(T(10), v2.capacity());
+
+ // test when insertion range fits inside capacity
+ // and the size of the insertion is equal to the number
+ // of displaced elements
+ Vector v3(5);
+ v3.reserve(10);
+ thrust::sequence(v3.begin(), v3.end());
+
+ new_size = v3.size() + v1.size();
+ insertion_size = v1.end() - v1.begin();
+ num_displaced = v3.end() - v3.begin();
+
+ ASSERT_EQ(true, v3.capacity() >= new_size);
+ ASSERT_EQ(true, insertion_size == num_displaced);
+
+ v3.insert(v3.begin(),
+ v1.begin(), v1.end());
+
+ ASSERT_EQ(T(0), v3[0]);
+ ASSERT_EQ(T(1), v3[1]);
+ ASSERT_EQ(T(2), v3[2]);
+ ASSERT_EQ(T(3), v3[3]);
+ ASSERT_EQ(T(4), v3[4]);
+
+ ASSERT_EQ(T(0), v3[5]);
+ ASSERT_EQ(T(1), v3[6]);
+ ASSERT_EQ(T(2), v3[7]);
+ ASSERT_EQ(T(3), v3[8]);
+ ASSERT_EQ(T(4), v3[9]);
+
+ ASSERT_EQ(10, v3.size());
+ ASSERT_EQ(10, v3.capacity());
+
+ // test when insertion range fits inside capacity
+ // and the size of the insertion is less than the
+ // number of displaced elements
+ Vector v4(5);
+ v4.reserve(10);
+ thrust::sequence(v4.begin(), v4.end());
+
+ new_size = v4.size() + v1.size();
+ insertion_size = (v1.begin() + 3) - v1.begin();
+ num_displaced = v4.end() - (v4.begin() + 1);
+
+ ASSERT_EQ(true, v4.capacity() >= new_size);
+ ASSERT_EQ(true, insertion_size < num_displaced);
+
+ v4.insert(v4.begin() + 1,
+ v1.begin(), v1.begin() + 3);
+
+ ASSERT_EQ(T(0), v4[0]);
+
+ ASSERT_EQ(T(0), v4[1]);
+ ASSERT_EQ(T(1), v4[2]);
+ ASSERT_EQ(T(2), v4[3]);
+
+ ASSERT_EQ(T(1), v4[4]);
+ ASSERT_EQ(T(2), v4[5]);
+ ASSERT_EQ(T(3), v4[6]);
+ ASSERT_EQ(T(4), v4[7]);
+
+ ASSERT_EQ(8, v4.size());
+ ASSERT_EQ(10, v4.capacity());
+
+ // test when insertion range does not fit inside capacity
+ Vector v5(5);
+ thrust::sequence(v5.begin(), v5.end());
+
+ new_size = v5.size() + v1.size();
+
+ ASSERT_EQ(true, v5.capacity() < new_size);
+
+ v5.insert(v5.begin() + 1,
+ v1.begin(), v1.end());
+
+ ASSERT_EQ(T(0), v5[0]);
+
+ ASSERT_EQ(T(0), v5[1]);
+ ASSERT_EQ(T(1), v5[2]);
+ ASSERT_EQ(T(2), v5[3]);
+ ASSERT_EQ(T(3), v5[4]);
+ ASSERT_EQ(T(4), v5[5]);
+
+ ASSERT_EQ(T(1), v5[6]);
+ ASSERT_EQ(T(2), v5[7]);
+ ASSERT_EQ(T(3), v5[8]);
+ ASSERT_EQ(T(4), v5[9]);
+
+ ASSERT_EQ(10, v5.size());
+}
+
+TYPED_TEST(VectorInsertTests, TestVectorRangeInsert)
+{
+ 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_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,
+ std::numeric_limits<T>::min(),
+ 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,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max()
+ );
+ }
+
+ thrust::device_vector<T> d_src = h_src;
+ thrust::device_vector<T> d_dst = h_dst;
+
+ // choose insertion range at random
+ size_t begin = size > 0 ? (size_t)h_src[size] % size : 0;
+ size_t end = size > 0 ? (size_t)h_src[size+1] % size : 0;
+ if(end < begin) thrust::swap(begin,end);
+
+ // choose insertion position at random
+ size_t position = size > 0 ? (size_t)h_src[size+2] % size : 0;
+
+ // insert on host
+ h_dst.insert(h_dst.begin() + position,
+ h_src.begin() + begin,
+ h_src.begin() + end);
+
+ // insert on device
+ d_dst.insert(d_dst.begin() + position,
+ d_src.begin() + begin,
+ d_src.begin() + end);
+
+ ASSERT_EQ(h_dst, d_dst);
+ }
+}
+
+TYPED_TEST(VectorInsertTests, TestVectorFillInsertSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector v1(3);
+ v1.reserve(10);
+ thrust::sequence(v1.begin(), v1.end());
+
+ size_t insertion_size = 5;
+ size_t new_size = v1.size() + insertion_size;
+ size_t num_displaced = v1.end() - (v1.begin() + 1);
+
+ ASSERT_EQ(true, v1.capacity() >= new_size);
+ ASSERT_EQ(true, insertion_size > num_displaced);
+
+ v1.insert(v1.begin() + 1, insertion_size, 13);
+
+ ASSERT_EQ(T(0), v1[0]);
+
+ ASSERT_EQ(T(13), v1[1]);
+ ASSERT_EQ(T(13), v1[2]);
+ ASSERT_EQ(T(13), v1[3]);
+ ASSERT_EQ(T(13), v1[4]);
+ ASSERT_EQ(T(13), v1[5]);
+
+ ASSERT_EQ(T(1), v1[6]);
+ ASSERT_EQ(T(2), v1[7]);
+
+ ASSERT_EQ(8, v1.size());
+ ASSERT_EQ(10, v1.capacity());
+
+ // test when insertion range fits inside capacity
+ // and the size of the insertion is equal to the number
+ // of displaced elements
+ Vector v2(5);
+ v2.reserve(10);
+ thrust::sequence(v2.begin(), v2.end());
+
+ insertion_size = 5;
+ new_size = v2.size() + insertion_size;
+ num_displaced = v2.end() - v2.begin();
+
+ ASSERT_EQ(true, v2.capacity() >= new_size);
+ ASSERT_EQ(true, insertion_size == num_displaced);
+
+ v2.insert(v2.begin(), insertion_size, 13);
+
+ ASSERT_EQ(T(13), v2[0]);
+ ASSERT_EQ(T(13), v2[1]);
+ ASSERT_EQ(T(13), v2[2]);
+ ASSERT_EQ(T(13), v2[3]);
+ ASSERT_EQ(T(13), v2[4]);
+
+ ASSERT_EQ(T(0), v2[5]);
+ ASSERT_EQ(T(1), v2[6]);
+ ASSERT_EQ(T(2), v2[7]);
+ ASSERT_EQ(T(3), v2[8]);
+ ASSERT_EQ(T(4), v2[9]);
+
+ ASSERT_EQ(10, v2.size());
+ ASSERT_EQ(10, v2.capacity());
+
+ // test when insertion range fits inside capacity
+ // and the size of the insertion is less than the
+ // number of displaced elements
+ Vector v3(5);
+ v3.reserve(10);
+ thrust::sequence(v3.begin(), v3.end());
+
+ insertion_size = 3;
+ new_size = v3.size() + insertion_size;
+ num_displaced = v3.end() - (v3.begin() + 1);
+
+ ASSERT_EQ(true, v3.capacity() >= new_size);
+ ASSERT_EQ(true, insertion_size < num_displaced);
+
+ v3.insert(v3.begin() + 1, insertion_size, 13);
+
+ ASSERT_EQ(T(0), v3[0]);
+
+ ASSERT_EQ(T(13), v3[1]);
+ ASSERT_EQ(T(13), v3[2]);
+ ASSERT_EQ(T(13), v3[3]);
+
+ ASSERT_EQ(T(1), v3[4]);
+ ASSERT_EQ(T(2), v3[5]);
+ ASSERT_EQ(T(3), v3[6]);
+ ASSERT_EQ(T(4), v3[7]);
+
+ ASSERT_EQ(8, v3.size());
+ ASSERT_EQ(10, v3.capacity());
+
+ // test when insertion range does not fit inside capacity
+ Vector v4(5);
+ thrust::sequence(v4.begin(), v4.end());
+
+ insertion_size = 5;
+ new_size = v4.size() + insertion_size;
+
+ ASSERT_EQ(true, v4.capacity() < new_size);
+
+ v4.insert(v4.begin() + 1, insertion_size, 13);
+
+ ASSERT_EQ(T(0), v4[0]);
+
+ ASSERT_EQ(T(13), v4[1]);
+ ASSERT_EQ(T(13), v4[2]);
+ ASSERT_EQ(T(13), v4[3]);
+ ASSERT_EQ(T(13), v4[4]);
+ ASSERT_EQ(T(13), v4[5]);
+
+ ASSERT_EQ(T(1), v4[6]);
+ ASSERT_EQ(T(2), v4[7]);
+ ASSERT_EQ(T(3), v4[8]);
+ ASSERT_EQ(T(4), v4[9]);
+
+ ASSERT_EQ(10, v4.size());
+}
+
+TYPED_TEST(VectorInsertTests, TestVectorFillInsert)
+{
+ 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_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,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max()
+ );
+ }
+
+ thrust::device_vector<T> d_dst = h_dst;
+
+ // choose insertion position at random
+ size_t position = size > 0 ? (size_t)h_dst[size] % size : 0;
+
+ // choose insertion size at random
+ size_t insertion_size = size > 0 ? (size_t)h_dst[size] % size : 13;
+
+ // insert on host
+ h_dst.insert(h_dst.begin() + position,
+ insertion_size,
+ T(13));
+
+ // insert on device
+ d_dst.insert(d_dst.begin() + position,
+ insertion_size,
+ T(13));
+
+ ASSERT_EQ(h_dst, d_dst);
+ }
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

Add vector insert and scatter tests