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
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
|
/*
* Copyright 2008-2013 NVIDIA Corporation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#pragma once
#include <thrust/detail/config.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/iterator_traits.h>
#include <thrust/detail/numeric_traits.h>
#include <thrust/detail/type_traits.h>
#include <cstddef>
THRUST_NAMESPACE_BEGIN
// forward declaration of counting_iterator
template <typename Incrementable, typename System, typename Traversal, typename Difference>
class counting_iterator;
namespace detail
{
template <typename Incrementable, typename System, typename Traversal, typename Difference>
struct counting_iterator_base
{
typedef typename thrust::detail::eval_if<
// use any_system_tag if we are given use_default
thrust::detail::is_same<System,use_default>::value,
thrust::detail::identity_<thrust::any_system_tag>,
thrust::detail::identity_<System>
>::type system;
typedef typename thrust::detail::ia_dflt_help<
Traversal,
thrust::detail::eval_if<
thrust::detail::is_numeric<Incrementable>::value,
thrust::detail::identity_<random_access_traversal_tag>,
thrust::iterator_traversal<Incrementable>
>
>::type traversal;
// unlike Boost, we explicitly use std::ptrdiff_t as the difference type
// for floating point counting_iterators
typedef typename thrust::detail::ia_dflt_help<
Difference,
thrust::detail::eval_if<
thrust::detail::is_numeric<Incrementable>::value,
thrust::detail::eval_if<
thrust::detail::is_integral<Incrementable>::value,
thrust::detail::numeric_difference<Incrementable>,
thrust::detail::identity_<std::ptrdiff_t>
>,
thrust::iterator_difference<Incrementable>
>
>::type difference;
// our implementation departs from Boost's in that counting_iterator::dereference
// returns a copy of its counter, rather than a reference to it. returning a reference
// to the internal state of an iterator causes subtle bugs (consider the temporary
// iterator created in the expression *(iter + i)) and has no compelling use case
typedef thrust::iterator_adaptor<
counting_iterator<Incrementable, System, Traversal, Difference>, // self
Incrementable, // Base
Incrementable, // XXX we may need to pass const here as Boost does
system,
traversal,
Incrementable,
difference
> type;
}; // end counting_iterator_base
template<typename Difference, typename Incrementable1, typename Incrementable2>
struct iterator_distance
{
__host__ __device__
static Difference distance(Incrementable1 x, Incrementable2 y)
{
return y - x;
}
};
template<typename Difference, typename Incrementable1, typename Incrementable2>
struct number_distance
{
__host__ __device__
static Difference distance(Incrementable1 x, Incrementable2 y)
{
return static_cast<Difference>(numeric_distance(x,y));
}
};
template<typename Difference, typename Incrementable1, typename Incrementable2, typename Enable = void>
struct counting_iterator_equal
{
__host__ __device__
static bool equal(Incrementable1 x, Incrementable2 y)
{
return x == y;
}
};
// specialization for floating point equality
template<typename Difference, typename Incrementable1, typename Incrementable2>
struct counting_iterator_equal<
Difference,
Incrementable1,
Incrementable2,
typename thrust::detail::enable_if<
thrust::detail::is_floating_point<Incrementable1>::value ||
thrust::detail::is_floating_point<Incrementable2>::value
>::type
>
{
__host__ __device__
static bool equal(Incrementable1 x, Incrementable2 y)
{
typedef number_distance<Difference,Incrementable1,Incrementable2> d;
return d::distance(x,y) == 0;
}
};
} // end detail
THRUST_NAMESPACE_END
|
