forked from NVIDIA/MinkowskiEngine
-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy pathcoordinate_map_functors.cuh
227 lines (182 loc) · 7.32 KB
/
coordinate_map_functors.cuh
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
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
/*
* Copyright 2020 NVIDIA CORPORATION.
*
* 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.
*/
#ifndef COORDINATE_FUNCTORS_CUH
#define COORDINATE_FUNCTORS_CUH
#include "coordinate.hpp"
#include "types.hpp"
#include <thrust/functional.h>
#include <thrust/host_vector.h>
#include <thrust/pair.h>
namespace minkowski {
namespace detail {
template <typename Key, typename Element, typename Equality> struct is_used {
using value_type = thrust::pair<Key, Element>;
is_used(Key &&unused, Equality &&equal)
: m_unused_key(unused), m_equal(equal) {}
is_used(Key const &unused, Equality const &equal)
: m_unused_key(unused), m_equal(equal) {}
__host__ __device__ bool operator()(value_type const &x) {
return !m_equal(x.first, m_unused_key);
}
Key const m_unused_key;
Equality const m_equal;
};
template <typename coordinate_type, typename map_type> struct get_element {
using value_type = thrust::pair<coordinate<coordinate_type>, uint32_t>;
get_element(map_type &map) : m_map(map) {}
__device__ uint32_t operator()(value_type const &x) { return x.second; }
map_type const m_map;
};
template <typename coordinate_type, typename map_type, typename mapped_iterator>
struct insert_coordinate {
using value_type = typename map_type::value_type;
using mapped_type = typename map_type::mapped_type;
/**
* insert_coordinate functor constructor
* @param map
* @param p_coordinate a pointer to the start of the coordinate
* @param value_iter a mapped_iterator that points to the begin. This could be
* a pointer or an iterator that supports operat+(int) and operator*().
* @param coordinate_size
*/
insert_coordinate(map_type &map, // underlying map
coordinate_type const *p_coordinate, // key coordinate begin
mapped_iterator row_iter, mapped_iterator map_iter,
uint32_t const coordinate_size) // coordinate size
: m_coordinate_size{coordinate_size}, m_coordinates{p_coordinate},
m_row_iter{row_iter}, m_map_iter{map_iter}, m_map{map} {}
/*
* @brief insert a <coordinate, row index> pair into the unordered_map
*
* @return thrust::pair<bool, uint32_t> of a success flag and the current
* index.
*/
__device__ bool operator()(uint32_t i) {
auto coord =
coordinate<coordinate_type>{&m_coordinates[i * m_coordinate_size]};
value_type pair = thrust::make_pair(coord, i);
// Returns pair<iterator, (bool)insert_success>
auto result = m_map.insert(pair);
m_row_iter[i] = i;
m_map_iter[i] = result.first.offset();
return result.second;
}
size_t const m_coordinate_size;
coordinate_type const *m_coordinates;
mapped_iterator m_row_iter;
mapped_iterator m_map_iter;
map_type m_map;
};
template <typename coordinate_type, typename map_type>
struct find_coordinate
: public thrust::unary_function<uint32_t,
thrust::pair<uint32_t, uint32_t>> {
using mapped_type = typename map_type::mapped_type;
using return_type = thrust::pair<mapped_type, mapped_type>;
find_coordinate(map_type const &_map, coordinate_type const *_d_ptr,
mapped_type const unused_element, size_t const _size)
: m_coordinate_size{_size},
m_unused_element(unused_element), m_coordinates{_d_ptr}, m_map{_map} {}
__device__ mapped_type operator()(uint32_t i) {
auto coord =
coordinate<coordinate_type>{&m_coordinates[i * m_coordinate_size]};
auto result = m_map.find(coord);
if (result == m_map.end()) {
return m_unused_element;
}
return result->second;
}
size_t const m_coordinate_size;
mapped_type const m_unused_element;
coordinate_type const *m_coordinates;
map_type const m_map;
};
template <typename coordinate_type, typename map_type> struct update_value {
using mapped_type = typename map_type::mapped_type;
update_value(map_type &_map, coordinate_type const *coordinates,
uint32_t const *valid_index, size_t const _size)
: m_coordinate_size{_size}, m_coordinates{coordinates},
m_valid_index(valid_index), m_map{_map} {}
__device__ void operator()(uint32_t i) {
auto coord = coordinate<coordinate_type>{
&m_coordinates[m_valid_index[i] * m_coordinate_size]};
auto result = m_map.find(coord);
result->second = i;
}
size_t const m_coordinate_size;
coordinate_type const *m_coordinates;
uint32_t const *m_valid_index;
map_type m_map;
};
template <typename index_type, typename map_type>
struct update_value_with_offset {
update_value_with_offset(map_type &_map, index_type const *valid_map_offset)
: m_valid_map_offset(valid_map_offset), m_map{_map} {}
__device__ void operator()(index_type i) {
auto &result = m_map.data()[m_valid_map_offset[i]];
result.second = i;
}
index_type const *m_valid_map_offset;
map_type m_map;
};
template <typename T> struct is_first {
is_first(T value) : m_value(value) {}
template <typename Tuple>
inline __device__ bool operator()(Tuple const &item) const {
return thrust::get<0>(item) == m_value;
}
T m_value;
};
template <typename coordinate_type, typename mapped_type>
struct is_unused_pair {
is_unused_pair(mapped_type const unused_element)
: m_unused_element(unused_element) {}
template <typename Tuple>
inline __device__ bool operator()(Tuple const &item) const {
return thrust::get<1>(item) == m_unused_element;
}
mapped_type const m_unused_element;
};
template <typename T, typename pair_type, typename pair_iterator>
struct split_functor {
split_functor(pair_iterator begin, T *firsts, T *seconds)
: m_begin(begin), m_firsts(firsts), m_seconds(seconds) {}
__device__ void operator()(uint32_t i) {
pair_type const item = *(m_begin + i);
m_firsts[i] = item.first;
m_seconds[i] = item.second;
}
pair_iterator m_begin;
T *m_firsts;
T *m_seconds;
};
template <typename T> struct min_size_functor {
using min_size_type = thrust::tuple<T, T>;
__host__ __device__ min_size_type operator()(const min_size_type &lhs,
const min_size_type &rhs) {
return thrust::make_tuple(min(thrust::get<0>(lhs), thrust::get<0>(rhs)),
thrust::get<1>(lhs) + thrust::get<1>(rhs));
}
};
} // end namespace detail
} // end namespace minkowski
#endif // COORDS_FUNCTORS_CUH