Skip to content

HTTPS clone URL

Subversion checkout URL

You can clone with
or
.
Download ZIP
Fetching contributors…

Cannot retrieve contributors at this time

341 lines (252 sloc) 10.365 kb
#include <unittest/unittest.h>
#include <thrust/scatter.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/sequence.h>
#include <thrust/fill.h>
struct my_system : thrust::device_system<my_system> {};
template <class Vector>
void TestScatterSimple(void)
{
typedef typename Vector::value_type T;
Vector map(5); // scatter indices
Vector src(5); // source vector
Vector dst(8); // destination vector
map[0] = 6; map[1] = 3; map[2] = 1; map[3] = 7; map[4] = 2;
src[0] = 0; src[1] = 1; src[2] = 2; src[3] = 3; src[4] = 4;
dst[0] = 0; dst[1] = 0; dst[2] = 0; dst[3] = 0; dst[4] = 0; dst[5] = 0; dst[6] = 0; dst[7] = 0;
thrust::scatter(src.begin(), src.end(), map.begin(), dst.begin());
ASSERT_EQUAL(dst[0], 0);
ASSERT_EQUAL(dst[1], 2);
ASSERT_EQUAL(dst[2], 4);
ASSERT_EQUAL(dst[3], 1);
ASSERT_EQUAL(dst[4], 0);
ASSERT_EQUAL(dst[5], 0);
ASSERT_EQUAL(dst[6], 0);
ASSERT_EQUAL(dst[7], 3);
}
DECLARE_VECTOR_UNITTEST(TestScatterSimple);
template<typename InputIterator1,
typename InputIterator2,
typename RandomAccessIterator>
void scatter(my_system,
InputIterator1,
InputIterator1,
InputIterator2,
RandomAccessIterator output)
{
*output = 13;
}
void TestScatterDispatchExplicit()
{
thrust::device_vector<int> vec(1);
my_system sys;
thrust::scatter(sys,
vec.begin(),
vec.begin(),
vec.begin(),
vec.begin());
ASSERT_EQUAL(13, vec.front());
}
DECLARE_UNITTEST(TestScatterDispatchExplicit);
void TestScatterDispatchImplicit()
{
thrust::device_vector<int> vec(1);
thrust::scatter(thrust::retag<my_system>(vec.begin()),
thrust::retag<my_system>(vec.begin()),
thrust::retag<my_system>(vec.begin()),
thrust::retag<my_system>(vec.begin()));
ASSERT_EQUAL(13, vec.front());
}
DECLARE_UNITTEST(TestScatterDispatchImplicit);
template <typename T>
void TestScatter(const size_t n)
{
const size_t output_size = std::min((size_t) 10, 2 * n);
thrust::host_vector<T> h_input(n, (T) 1);
thrust::device_vector<T> d_input(n, (T) 1);
thrust::host_vector<unsigned int> h_map = unittest::random_integers<unsigned int>(n);
for(size_t i = 0; i < n; 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_EQUAL(h_output, d_output);
}
DECLARE_VARIABLE_UNITTEST(TestScatter);
template <typename T>
void TestScatterToDiscardIterator(const size_t n)
{
const size_t output_size = std::min((size_t) 10, 2 * n);
thrust::host_vector<T> h_input(n, (T) 1);
thrust::device_vector<T> d_input(n, (T) 1);
thrust::host_vector<unsigned int> h_map = unittest::random_integers<unsigned int>(n);
for(size_t i = 0; i < n; 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
}
DECLARE_VARIABLE_UNITTEST(TestScatterToDiscardIterator);
template <class Vector>
void TestScatterIfSimple(void)
{
typedef typename Vector::value_type T;
Vector flg(5); // predicate array
Vector map(5); // scatter indices
Vector src(5); // source vector
Vector dst(8); // destination vector
flg[0] = 0; flg[1] = 1; flg[2] = 0; flg[3] = 1; flg[4] = 0;
map[0] = 6; map[1] = 3; map[2] = 1; map[3] = 7; map[4] = 2;
src[0] = 0; src[1] = 1; src[2] = 2; src[3] = 3; src[4] = 4;
dst[0] = 0; dst[1] = 0; dst[2] = 0; dst[3] = 0; dst[4] = 0; dst[5] = 0; dst[6] = 0; dst[7] = 0;
thrust::scatter_if(src.begin(), src.end(), map.begin(), flg.begin(), dst.begin());
ASSERT_EQUAL(dst[0], 0);
ASSERT_EQUAL(dst[1], 0);
ASSERT_EQUAL(dst[2], 0);
ASSERT_EQUAL(dst[3], 1);
ASSERT_EQUAL(dst[4], 0);
ASSERT_EQUAL(dst[5], 0);
ASSERT_EQUAL(dst[6], 0);
ASSERT_EQUAL(dst[7], 3);
}
DECLARE_VECTOR_UNITTEST(TestScatterIfSimple);
template<typename InputIterator1,
typename InputIterator2,
typename InputIterator3,
typename RandomAccessIterator>
void scatter_if(my_system,
InputIterator1,
InputIterator1,
InputIterator2,
InputIterator3,
RandomAccessIterator output)
{
*output = 13;
}
void TestScatterIfDispatchExplicit()
{
thrust::device_vector<int> vec(1);
my_system sys;
thrust::scatter_if(sys,
vec.begin(),
vec.begin(),
vec.begin(),
vec.begin(),
vec.begin());
ASSERT_EQUAL(13, vec.front());
}
DECLARE_UNITTEST(TestScatterIfDispatchExplicit);
void TestScatterIfDispatchImplicit()
{
thrust::device_vector<int> vec(1);
thrust::scatter_if(thrust::retag<my_system>(vec.begin()),
thrust::retag<my_system>(vec.begin()),
thrust::retag<my_system>(vec.begin()),
thrust::retag<my_system>(vec.begin()),
thrust::retag<my_system>(vec.begin()));
ASSERT_EQUAL(13, vec.front());
}
DECLARE_UNITTEST(TestScatterIfDispatchImplicit);
template <typename T>
class is_even_scatter_if
{
public:
__host__ __device__ bool operator()(const T i) const { return (i % 2) == 0; }
};
template <typename T>
void TestScatterIf(const size_t n)
{
const size_t output_size = std::min((size_t) 10, 2 * n);
thrust::host_vector<T> h_input(n, (T) 1);
thrust::device_vector<T> d_input(n, (T) 1);
thrust::host_vector<unsigned int> h_map = unittest::random_integers<unsigned int>(n);
for(size_t i = 0; i < n; 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_EQUAL(h_output, d_output);
}
DECLARE_VARIABLE_UNITTEST(TestScatterIf);
template <typename T>
void TestScatterIfToDiscardIterator(const size_t n)
{
const size_t output_size = std::min((size_t) 10, 2 * n);
thrust::host_vector<T> h_input(n, (T) 1);
thrust::device_vector<T> d_input(n, (T) 1);
thrust::host_vector<unsigned int> h_map = unittest::random_integers<unsigned int>(n);
for(size_t i = 0; i < n; 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>());
}
DECLARE_VARIABLE_UNITTEST(TestScatterIfToDiscardIterator);
template <typename Vector>
void TestScatterCountingIterator(void)
{
typedef typename Vector::value_type T;
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_EQUAL(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_EQUAL(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_EQUAL(output, map);
}
DECLARE_VECTOR_UNITTEST(TestScatterCountingIterator);
template <typename Vector>
void TestScatterIfCountingIterator(void)
{
typedef typename Vector::value_type T;
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_EQUAL(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_EQUAL(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_EQUAL(output, map);
}
DECLARE_VECTOR_UNITTEST(TestScatterIfCountingIterator);
Jump to Line
Something went wrong with that request. Please try again.