Extend scatter tests and added fill, allocator and device_delete tests
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -56,9 +56,12 @@ endfunction()
# Tests
# ****************************************************************************
+add_hipthrust_test("thrust.hip.allocator" test_allocator.cpp)
add_hipthrust_test("thrust.hip.copy" test_copy.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.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
test/test_allocator.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/device_malloc_allocator.h>
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+#include <thrust/system/cpp/vector.h>
+#include <memory>
+
+// 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
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+struct my_allocator_with_custom_construct1
+ : thrust::device_malloc_allocator<int>
+{
+ __host__ __device__
+ my_allocator_with_custom_construct1()
+ {}
+
+ template<typename T>
+ __host__ __device__
+ void construct(T *p)
+ {
+ *p = 13;
+ }
+};
+
+TEST(AllocatorTests, TestAllocatorCustomDefaultConstruct)
+{
+ thrust::device_vector<int> ref(10,13);
+ thrust::device_vector<int, my_allocator_with_custom_construct1> vec(10);
+
+ //ASSERT_EQ(ref, vec);
+}
+
+struct my_allocator_with_custom_construct2
+ : thrust::device_malloc_allocator<int>
+{
+ __host__ __device__
+ my_allocator_with_custom_construct2()
+ {}
+
+ template<typename T, typename Arg>
+ __host__ __device__
+ void construct(T *p, const Arg &)
+ {
+ *p = 13;
+ }
+};
+
+TEST(AllocatorTests, TestAllocatorCustomCopyConstruct)
+{
+ thrust::device_vector<int> ref(10,13);
+ thrust::device_vector<int> copy_from(10,7);
+ thrust::device_vector<int, my_allocator_with_custom_construct2> vec(copy_from.begin(), copy_from.end());
+
+ //ASSERT_EQ(ref, vec);
+}
+
+/*static int g_state;
+
+struct my_allocator_with_custom_destroy
+{
+ typedef int value_type;
+ typedef int & reference;
+ typedef const int & const_reference;
+
+ __host__
+ my_allocator_with_custom_destroy(){}
+
+ __host__
+ my_allocator_with_custom_destroy(const my_allocator_with_custom_destroy &other)
+ : use_me_to_alloc(other.use_me_to_alloc)
+ {}
+
+ __host__
+ ~my_allocator_with_custom_destroy(){}
+
+ template<typename T>
+ __host__ __device__
+ void destroy(T *p)
+ {
+ g_state = 13;
+ }
+
+ value_type *allocate(std::ptrdiff_t n)
+ {
+ return use_me_to_alloc.allocate(n);
+ }
+
+ void deallocate(value_type *ptr, std::ptrdiff_t n)
+ {
+ use_me_to_alloc.deallocate(ptr,n);
+ }
+
+ // use composition rather than inheritance
+ // to avoid inheriting std::allocator's member
+ // function construct
+ std::allocator<int> use_me_to_alloc;
+};
+
+TEST(AllocatorTests, TestAllocatorCustomDestroy)
+{
+ thrust::cpp::vector<int, my_allocator_with_custom_destroy> vec(10);
+
+ // destroy everything
+ vec.shrink_to_fit();
+
+ ASSERT_EQ(13, g_state);
+}*/
+
+/*struct my_minimal_allocator
+{
+ typedef int value_type;
+
+ // XXX ideally, we shouldn't require
+ // these two typedefs
+ typedef int & reference;
+ typedef const int & const_reference;
+
+ __host__
+ my_minimal_allocator(){}
+
+ __host__
+ my_minimal_allocator(const my_minimal_allocator &other)
+ : use_me_to_alloc(other.use_me_to_alloc)
+ {}
+
+ __host__
+ ~my_minimal_allocator(){}
+
+ value_type *allocate(std::ptrdiff_t n)
+ {
+ return use_me_to_alloc.allocate(n);
+ }
+
+ void deallocate(value_type *ptr, std::ptrdiff_t n)
+ {
+ use_me_to_alloc.deallocate(ptr,n);
+ }
+
+ std::allocator<int> use_me_to_alloc;
+};
+
+TEST(AllocatorTests, TestAllocatorMinimal)
+{
+ thrust::cpp::vector<int, my_minimal_allocator> vec(10, 13);
+
+ // XXX copy to h_vec because ASSERT_EQUAL doesn't know about cpp::vector
+ thrust::host_vector<int> h_vec(vec.begin(), vec.end());
+ thrust::host_vector<int> ref(10, 13);
+
+ ASSERT_EQ(ref, h_vec);
+}*/
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

ADD
test/test_device_delete.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/device_vector.h>
+#include <thrust/device_ptr.h>
+#include <thrust/device_new.h>
+#include <thrust/device_delete.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
+
+struct Foo
+{
+ __host__ __device__
+ Foo(void)
+ :set_me_upon_destruction(0)
+ {}
+
+ bool *set_me_upon_destruction;
+};
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+TEST(DeviceDelete, TestDeviceDeleteDestructorInvocation)
+{
+ thrust::device_vector<bool> destructor_flag(1, false);
+
+ thrust::device_ptr<Foo> foo_ptr = thrust::device_new<Foo>();
+
+ Foo exemplar;
+ exemplar.set_me_upon_destruction = thrust::raw_pointer_cast(&destructor_flag[0]);
+ *foo_ptr = exemplar;
+
+ ASSERT_EQ(false, destructor_flag[0]);
+
+ //thrust::device_delete(foo_ptr);
+
+ //ASSERT_EQ(true, destructor_flag[0]);
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

ADD
test/test_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.
+
+#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 FillTests : 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>>
+> FillTestsParams;
+
+TYPED_TEST_CASE(FillTests, FillTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+TYPED_TEST(FillTests, TestFillSimple)
+{
+ 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);
+
+ thrust::fill(v.begin() + 1, v.begin() + 4, T(7));
+
+ ASSERT_EQ(v[0], T(0));
+ ASSERT_EQ(v[1], T(7));
+ ASSERT_EQ(v[2], T(7));
+ ASSERT_EQ(v[3], T(7));
+ ASSERT_EQ(v[4], T(4));
+
+ thrust::fill(v.begin() + 0, v.begin() + 3, T(8));
+
+ ASSERT_EQ(v[0], T(8));
+ ASSERT_EQ(v[1], T(8));
+ ASSERT_EQ(v[2], T(8));
+ ASSERT_EQ(v[3], T(7));
+ ASSERT_EQ(v[4], T(4));
+
+ thrust::fill(v.begin() + 2, v.end(), T(9));
+
+ ASSERT_EQ(v[0], T(8));
+ ASSERT_EQ(v[1], T(8));
+ ASSERT_EQ(v[2], T(9));
+ ASSERT_EQ(v[3], T(9));
+ ASSERT_EQ(v[4], T(9));
+
+ thrust::fill(v.begin(), v.end(), T(1));
+
+ ASSERT_EQ(v[0], T(1));
+ ASSERT_EQ(v[1], T(1));
+ ASSERT_EQ(v[2], T(1));
+ ASSERT_EQ(v[3], T(1));
+ ASSERT_EQ(v[4], T(1));
+
+}
+
+TEST(FillTests, TestFillDiscardIterator)
+{
+ // there's no result to check because fill returns void
+ thrust::fill(thrust::discard_iterator<thrust::host_system_tag>(),
+ thrust::discard_iterator<thrust::host_system_tag>(10),
+ 13);
+
+ thrust::fill(thrust::discard_iterator<thrust::device_system_tag>(),
+ thrust::discard_iterator<thrust::device_system_tag>(10),
+ 13);
+}
+
+TYPED_TEST(FillTests, TestFillMixedTypes)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector v(4);
+
+ thrust::fill(v.begin(), v.end(), (long) 10);
+
+ ASSERT_EQ(v[0], T(10));
+ ASSERT_EQ(v[1], T(10));
+ ASSERT_EQ(v[2], T(10));
+ ASSERT_EQ(v[3], T(10));
+
+ thrust::fill(v.begin(), v.end(), (float) 20);
+
+ ASSERT_EQ(v[0], T(20));
+ ASSERT_EQ(v[1], T(20));
+ ASSERT_EQ(v[2], T(20));
+ ASSERT_EQ(v[3], T(20));
+}
+
+TYPED_TEST(FillTests, TestFill)
+{
+ 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::fill(h_data.begin() + std::min((size_t)1, size), h_data.begin() + std::min((size_t)3, size), T(0));
+ thrust::fill(d_data.begin() + std::min((size_t)1, size), d_data.begin() + std::min((size_t)3, size), T(0));
+
+ ASSERT_EQ(h_data, d_data);
+
+ thrust::fill(h_data.begin() + std::min((size_t)117, size), h_data.begin() + std::min((size_t)367, size), T(1));
+ thrust::fill(d_data.begin() + std::min((size_t)117, size), d_data.begin() + std::min((size_t)367, size), T(1));
+
+ ASSERT_EQ(h_data, d_data);
+
+ thrust::fill(h_data.begin() + std::min((size_t)8, size), h_data.begin() + std::min((size_t)259, size), T(2));
+ thrust::fill(d_data.begin() + std::min((size_t)8, size), d_data.begin() + std::min((size_t)259, size), T(2));
+
+ ASSERT_EQ(h_data, d_data);
+
+ thrust::fill(h_data.begin() + std::min((size_t)3, size), h_data.end(), T(3));
+ thrust::fill(d_data.begin() + std::min((size_t)3, size), d_data.end(), T(3));
+
+ ASSERT_EQ(h_data, d_data);
+
+ thrust::fill(h_data.begin(), h_data.end(), T(4));
+ thrust::fill(d_data.begin(), d_data.end(), T(4));
+
+ ASSERT_EQ(h_data, d_data);
+ }
+}
+
+TYPED_TEST(FillTests, TestFillNSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+ using Iterator = typename Vector::iterator;
+
+ Vector v(5);
+ v[0] = T(0); v[1] = T(1); v[2] = T(2); v[3] = T(3); v[4] = T(4);
+
+ Iterator iter = thrust::fill_n(v.begin() + 1, 3, T(7));
+
+ ASSERT_EQ(v[0], T(0));
+ ASSERT_EQ(v[1], T(7));
+ ASSERT_EQ(v[2], T(7));
+ ASSERT_EQ(v[3], T(7));
+ ASSERT_EQ(v[4], T(4));
+ ASSERT_EQ(v.begin() + 4, iter);
+
+ iter = thrust::fill_n(v.begin() + 0, 3, T(8));
+
+ ASSERT_EQ(v[0], T(8));
+ ASSERT_EQ(v[1], T(8));
+ ASSERT_EQ(v[2], T(8));
+ ASSERT_EQ(v[3], T(7));
+ ASSERT_EQ(v[4], T(4));
+ ASSERT_EQ(v.begin() + 3, iter);
+
+ iter = thrust::fill_n(v.begin() + 2, 3, T(9));
+
+ ASSERT_EQ(v[0], T(8));
+ ASSERT_EQ(v[1], T(8));
+ ASSERT_EQ(v[2], T(9));
+ ASSERT_EQ(v[3], T(9));
+ ASSERT_EQ(v[4], T(9));
+ ASSERT_EQ(v.end(), iter);
+
+ iter = thrust::fill_n(v.begin(), v.size(), T(1));
+
+ ASSERT_EQ(v[0], T(1));
+ ASSERT_EQ(v[1], T(1));
+ ASSERT_EQ(v[2], T(1));
+ ASSERT_EQ(v[3], T(1));
+ ASSERT_EQ(v[4], T(1));
+ ASSERT_EQ(v.end(), iter);
+
+}
+
+TEST(FillTests, TestFillNDiscardIterator)
+{
+ thrust::discard_iterator<thrust::host_system_tag> h_result =
+ thrust::fill_n(thrust::discard_iterator<thrust::host_system_tag>(),
+ 10,
+ 13);
+
+ thrust::discard_iterator<thrust::device_system_tag> d_result =
+ thrust::fill_n(thrust::discard_iterator<thrust::device_system_tag>(),
+ 10,
+ 13);
+
+ thrust::discard_iterator<> reference(10);
+
+ ASSERT_EQ(reference, h_result);
+ ASSERT_EQ(reference, d_result);
+}
+
+TYPED_TEST(FillTests, TestFillNMixedTypes)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+ using Iterator = typename Vector::iterator;
+
+ Vector v(4);
+
+ Iterator iter = thrust::fill_n(v.begin(), v.size(), (long) 10);
+
+ ASSERT_EQ(v[0], T(10));
+ ASSERT_EQ(v[1], T(10));
+ ASSERT_EQ(v[2], T(10));
+ ASSERT_EQ(v[3], T(10));
+ ASSERT_EQ(v.end(), iter);
+
+ iter = thrust::fill_n(v.begin(), v.size(), (float) 20);
+
+ ASSERT_EQ(v[0], T(20));
+ ASSERT_EQ(v[1], T(20));
+ ASSERT_EQ(v[2], T(20));
+ ASSERT_EQ(v[3], T(20));
+ ASSERT_EQ(v.end(), iter);
+}
+
+
+TYPED_TEST(FillTests, TestFillN)
+{
+ 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 begin_offset = std::min<size_t>(1,size);
+ thrust::fill_n(h_data.begin() + begin_offset, std::min((size_t)3, size) - begin_offset, T(0));
+ thrust::fill_n(d_data.begin() + begin_offset, std::min((size_t)3, size) - begin_offset, T(0));
+
+ ASSERT_EQ(h_data, d_data);
+
+ begin_offset = std::min<size_t>(117, size);
+ thrust::fill_n(h_data.begin() + begin_offset, std::min((size_t)367, size) - begin_offset, T(1));
+ thrust::fill_n(d_data.begin() + begin_offset, std::min((size_t)367, size) - begin_offset, T(1));
+
+ ASSERT_EQ(h_data, d_data);
+
+ begin_offset = std::min<size_t>(8, size);
+ thrust::fill_n(h_data.begin() + begin_offset, std::min((size_t)259, size) - begin_offset, T(2));
+ thrust::fill_n(d_data.begin() + begin_offset, std::min((size_t)259, size) - begin_offset, T(2));
+
+ ASSERT_EQ(h_data, d_data);
+
+ begin_offset = std::min<size_t>(3, size);
+ thrust::fill_n(h_data.begin() + begin_offset, h_data.size() - begin_offset, T(3));
+ thrust::fill_n(d_data.begin() + begin_offset, d_data.size() - begin_offset, T(3));
+
+ ASSERT_EQ(h_data, d_data);
+
+ thrust::fill_n(h_data.begin(), h_data.size(), T(4));
+ thrust::fill_n(d_data.begin(), d_data.size(), T(4));
+
+ ASSERT_EQ(h_data, d_data);
+ }
+}
+
+TYPED_TEST(FillTests, TestFillZipIterator)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector v1(3,T(0));
+ Vector v2(3,T(0));
+ Vector v3(3,T(0));
+
+ thrust::fill(thrust::make_zip_iterator(thrust::make_tuple(v1.begin(),v2.begin(),v3.begin())),
+ thrust::make_zip_iterator(thrust::make_tuple(v1.end(),v2.end(),v3.end())),
+ thrust::tuple<T,T,T>(4,7,13));
+
+ ASSERT_EQ(T(4), v1[0]);
+ ASSERT_EQ(T(4), v1[1]);
+ ASSERT_EQ(T(4), v1[2]);
+ ASSERT_EQ(T(7), v2[0]);
+ ASSERT_EQ(T(7), v2[1]);
+ ASSERT_EQ(T(7), v2[2]);
+ ASSERT_EQ(T(13), v3[0]);
+ ASSERT_EQ(T(13), v3[1]);
+ ASSERT_EQ(T(13), v3[2]);
+}
+
+// TODO: Implement tuple.
+/*TYPED_TEST(FillTests, TestFillTuple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+ using Tuple = typename thrust::tuple<T,T>;
+
+ thrust::host_vector<Tuple> h(3, Tuple(0,0));
+ thrust::device_vector<Tuple> d(3, Tuple(0,0));
+
+ thrust::fill(h.begin(), h.end(), Tuple(4,7));
+ thrust::fill(d.begin(), d.end(), Tuple(4,7));
+
+ ASSERT_EQ(h, d);
+};*/
+
+struct TypeWithTrivialAssigment
+{
+ int x, y, z;
+};
+
+TEST(FillTests, TestFillWithTrivialAssignment)
+{
+ using T = TypeWithTrivialAssigment;
+
+ thrust::host_vector<T> h(1);
+ thrust::device_vector<T> d(1);
+
+ ASSERT_EQ(h[0].x, 0);
+ ASSERT_EQ(h[0].y, 0);
+ ASSERT_EQ(h[0].z, 0);
+ ASSERT_EQ(static_cast<T>(d[0]).x, 0);
+ ASSERT_EQ(static_cast<T>(d[0]).y, 0);
+ ASSERT_EQ(static_cast<T>(d[0]).z, 0);
+
+ T val;
+ val.x = 10;
+ val.y = 20;
+ val.z = -1;
+
+ thrust::fill(h.begin(), h.end(), val);
+ thrust::fill(d.begin(), d.end(), val);
+
+ ASSERT_EQ(h[0].x, 10);
+ ASSERT_EQ(h[0].y, 20);
+ ASSERT_EQ(h[0].z, -1);
+ ASSERT_EQ(static_cast<T>(d[0]).x, 10);
+ ASSERT_EQ(static_cast<T>(d[0]).y, 20);
+ ASSERT_EQ(static_cast<T>(d[0]).z, -1);
+}
+
+struct TypeWithNonTrivialAssigment
+{
+ int x, y, z;
+
+ __host__ __device__
+ TypeWithNonTrivialAssigment() : x(0), y(0), z(0) {}
+
+ __host__ __device__
+ TypeWithNonTrivialAssigment& operator=(const TypeWithNonTrivialAssigment& t)
+ {
+ x = t.x;
+ y = t.y;
+ z = t.x + t.y;
+ return *this;
+ }
+
+ __host__ __device__
+ bool operator==(const TypeWithNonTrivialAssigment& t) const
+ {
+ return x == t.x && y == t.y && z == t.z;
+ }
+};
+
+TEST(FillTests, TestFillWithNonTrivialAssignment)
+{
+ using T = TypeWithNonTrivialAssigment;
+
+ thrust::host_vector<T> h(1);
+ thrust::device_vector<T> d(1);
+
+ ASSERT_EQ(h[0].x, 0);
+ ASSERT_EQ(h[0].y, 0);
+ ASSERT_EQ(h[0].z, 0);
+ ASSERT_EQ(static_cast<T>(d[0]).x, 0);
+ ASSERT_EQ(static_cast<T>(d[0]).y, 0);
+ ASSERT_EQ(static_cast<T>(d[0]).z, 0);
+
+ T val;
+ val.x = 10;
+ val.y = 20;
+ val.z = -1;
+
+ thrust::fill(h.begin(), h.end(), val);
+ thrust::fill(d.begin(), d.end(), val);
+
+ ASSERT_EQ(h[0].x, 10);
+ ASSERT_EQ(h[0].y, 20);
+ ASSERT_EQ(h[0].z, 30);
+ ASSERT_EQ(static_cast<T>(d[0]).x, 10);
+ ASSERT_EQ(static_cast<T>(d[0]).y, 20);
+ ASSERT_EQ(static_cast<T>(d[0]).z, 30);
+}
+
+template<typename ForwardIterator, typename T>
+void fill(my_system &system, ForwardIterator, ForwardIterator, const T&)
+{
+ system.validate_dispatch();
+}
+
+TEST(FillTests, TestFillDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::fill(sys, vec.begin(), vec.end(), 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename ForwardIterator, typename T>
+void fill(my_tag, ForwardIterator first, ForwardIterator, const T&)
+{
+ *first = 13;
+}
+
+TEST(FillTests, TestFillDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::fill(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.end()),
+ 0);
+
+ ASSERT_EQ(13, vec.front());
+}
+
+template<typename OutputIterator, typename Size, typename T>
+OutputIterator fill_n(my_system &system, OutputIterator first, Size, const T&)
+{
+ system.validate_dispatch();
+ return first;
+}
+
+TEST(FillTests, TestFillNDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::fill_n(sys, vec.begin(), vec.size(), 0);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename OutputIterator, typename Size, typename T>
+OutputIterator fill_n(my_tag, OutputIterator first, Size, const T&)
+{
+ *first = 13;
+ return first;
+}
+
+TEST(FillTests, TestFillNDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::fill_n(thrust::retag<my_tag>(vec.begin()),
+ vec.size(),
+ 0);
+
+ ASSERT_EQ(13, vec.front());
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

MODIFY
test/test_scatter.cpp
test/test_scatter.cpp
@@ -276,4 +276,179 @@ TEST(ScatterTests, TestScatterIfDispatchExplicit)
ASSERT_EQ(true, sys.is_valid());
}
+template<typename InputIterator1,
+ typename InputIterator2,
+ typename InputIterator3,
+ typename RandomAccessIterator>
+void scatter_if(my_tag,
+ InputIterator1,
+ InputIterator1,
+ InputIterator2,
+ InputIterator3,
+ RandomAccessIterator output)
+{
+ *output = 13;
+}
+
+TEST(ScatterTests, TestScatterIfDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::scatter_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()),
+ thrust::retag<my_tag>(vec.begin()));
+
+ ASSERT_EQ(13, vec.front());
+}
+
+template <typename T>
+class is_even_scatter_if
+{
+ public:
+ __host__ __device__ bool operator()(const T i) const { return (i % 2) == 0; }
+};
+
+TYPED_TEST(ScatterTests, TestScatterIf)
+{
+ 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_if(h_input.begin(), h_input.end(), h_map.begin(), h_map.begin(), h_output.begin(), is_even_scatter_if<unsigned int>());
+ thrust::scatter_if(d_input.begin(), d_input.end(), d_map.begin(), d_map.begin(), d_output.begin(), is_even_scatter_if<unsigned int>());
+
+ ASSERT_EQ(h_output, d_output);
+ }
+}
+
+TYPED_TEST(ScatterTests, TestScatterIfToDiscardIterator)
+{
+ 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_if(h_input.begin(), h_input.end(), h_map.begin(), h_map.begin(), thrust::make_discard_iterator(), is_even_scatter_if<unsigned int>());
+ thrust::scatter_if(d_input.begin(), d_input.end(), d_map.begin(), d_map.begin(), thrust::make_discard_iterator(), is_even_scatter_if<unsigned int>());
+ }
+}
+
+// TODO: Implement reduce in system
+/*TYPED_TEST(ScatterTests, TestScatterCountingIterator)
+{
+ using Vector = typename TestFixture::input_type;
+
+ Vector source(10);
+ thrust::sequence(source.begin(), source.end(), 0);
+
+ Vector map(10);
+ thrust::sequence(map.begin(), map.end(), 0);
+
+ Vector output(10);
+
+ // source has any_system_tag
+ thrust::fill(output.begin(), output.end(), 0);
+ thrust::scatter(thrust::make_counting_iterator(0), thrust::make_counting_iterator(10),
+ map.begin(),
+ output.begin());
+
+ ASSERT_EQ(output, map);
+
+ // map has any_system_tag
+ thrust::fill(output.begin(), output.end(), 0);
+ thrust::scatter(source.begin(), source.end(),
+ thrust::make_counting_iterator(0),
+ output.begin());
+
+ ASSERT_EQ(output, map);
+
+ // source and map have any_system_tag
+ thrust::fill(output.begin(), output.end(), 0);
+ thrust::scatter(thrust::make_counting_iterator(0), thrust::make_counting_iterator(10),
+ thrust::make_counting_iterator(0),
+ output.begin());
+
+ ASSERT_EQ(output, map);
+}
+
+TYPED_TEST(ScatterTests, TestScatterIfCountingIterator)
+{
+ using Vector = typename TestFixture::input_type;
+
+ Vector source(10);
+ thrust::sequence(source.begin(), source.end(), 0);
+
+ Vector map(10);
+ thrust::sequence(map.begin(), map.end(), 0);
+
+ Vector stencil(10, 1);
+
+ Vector output(10);
+
+ // source has any_system_tag
+ thrust::fill(output.begin(), output.end(), 0);
+ thrust::scatter_if(thrust::make_counting_iterator(0), thrust::make_counting_iterator(10),
+ map.begin(),
+ stencil.begin(),
+ output.begin());
+
+ ASSERT_EQ(output, map);
+
+ // map has any_system_tag
+ thrust::fill(output.begin(), output.end(), 0);
+ thrust::scatter_if(source.begin(), source.end(),
+ thrust::make_counting_iterator(0),
+ stencil.begin(),
+ output.begin());
+
+ ASSERT_EQ(output, map);
+
+ // source and map have any_system_tag
+ thrust::fill(output.begin(), output.end(), 0);
+ thrust::scatter_if(thrust::make_counting_iterator(0), thrust::make_counting_iterator(10),
+ thrust::make_counting_iterator(0),
+ stencil.begin(),
+ output.begin());
+
+ ASSERT_EQ(output, map);
+}*/
+
#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

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

Extend scatter tests and added fill, allocator and device_delete tests