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