1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
|
#include <thrust/copy.h>
#include <thrust/device_vector.h>
#include <thrust/fill.h>
#include <thrust/functional.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <iostream>
#include "include/host_device.h"
// this example illustrates how to tile a range multiple times
// examples:
// tiled_range([0, 1, 2, 3], 1) -> [0, 1, 2, 3]
// tiled_range([0, 1, 2, 3], 2) -> [0, 1, 2, 3, 0, 1, 2, 3]
// tiled_range([0, 1, 2, 3], 3) -> [0, 1, 2, 3, 0, 1, 2, 3, 0, 1, 2, 3]
// ...
template <typename Iterator>
class tiled_range
{
public:
typedef typename thrust::iterator_difference<Iterator>::type difference_type;
struct tile_functor : public thrust::unary_function<difference_type, difference_type>
{
difference_type tile_size;
tile_functor(difference_type tile_size)
: tile_size(tile_size)
{}
__host__ __device__ difference_type operator()(const difference_type& i) const
{
return i % tile_size;
}
};
typedef typename thrust::counting_iterator<difference_type> CountingIterator;
typedef typename thrust::transform_iterator<tile_functor, CountingIterator> TransformIterator;
typedef typename thrust::permutation_iterator<Iterator, TransformIterator> PermutationIterator;
// type of the tiled_range iterator
typedef PermutationIterator iterator;
// construct repeated_range for the range [first,last)
tiled_range(Iterator first, Iterator last, difference_type tiles)
: first(first)
, last(last)
, tiles(tiles)
{}
iterator begin() const
{
return PermutationIterator(first, TransformIterator(CountingIterator(0), tile_functor(last - first)));
}
iterator end() const
{
return begin() + tiles * (last - first);
}
protected:
Iterator first;
Iterator last;
difference_type tiles;
};
int main()
{
thrust::device_vector<int> data(4);
data[0] = 10;
data[1] = 20;
data[2] = 30;
data[3] = 40;
// print the initial data
std::cout << "range ";
thrust::copy(data.begin(), data.end(), std::ostream_iterator<int>(std::cout, " "));
std::cout << std::endl;
typedef thrust::device_vector<int>::iterator Iterator;
// create tiled_range with two tiles
tiled_range<Iterator> two(data.begin(), data.end(), 2);
std::cout << "two tiles: ";
thrust::copy(two.begin(), two.end(), std::ostream_iterator<int>(std::cout, " "));
std::cout << std::endl;
// create tiled_range with three tiles
tiled_range<Iterator> three(data.begin(), data.end(), 3);
std::cout << "three tiles: ";
thrust::copy(three.begin(), three.end(), std::ostream_iterator<int>(std::cout, " "));
std::cout << std::endl;
return 0;
}
|
