|
21 | 21 | #include <rmm/exec_policy.hpp> |
22 | 22 |
|
23 | 23 | #include <thrust/copy.h> |
| 24 | +#include <thrust/remove.h> |
24 | 25 | #include <thrust/shuffle.h> |
25 | 26 | #include <thrust/sort.h> |
| 27 | +#include <thrust/unique.h> |
26 | 28 |
|
27 | 29 | namespace cugraph { |
28 | 30 | namespace test { |
@@ -164,5 +166,132 @@ template rmm::device_uvector<int32_t> randomly_select(raft::handle_t const& hand |
164 | 166 | template rmm::device_uvector<int64_t> randomly_select(raft::handle_t const& handle, |
165 | 167 | rmm::device_uvector<int64_t> const& input, |
166 | 168 | size_t count); |
| 169 | + |
| 170 | +template <typename vertex_t, typename weight_t> |
| 171 | +void remove_self_loops(raft::handle_t const& handle, |
| 172 | + rmm::device_uvector<vertex_t>& d_src_v /* [INOUT] */, |
| 173 | + rmm::device_uvector<vertex_t>& d_dst_v /* [INOUT] */, |
| 174 | + std::optional<rmm::device_uvector<weight_t>>& d_weight_v /* [INOUT] */) |
| 175 | +{ |
| 176 | + if (d_weight_v) { |
| 177 | + auto edge_first = thrust::make_zip_iterator( |
| 178 | + thrust::make_tuple(d_src_v.begin(), d_dst_v.begin(), (*d_weight_v).begin())); |
| 179 | + d_src_v.resize( |
| 180 | + thrust::distance(edge_first, |
| 181 | + thrust::remove_if( |
| 182 | + handle.get_thrust_policy(), |
| 183 | + edge_first, |
| 184 | + edge_first + d_src_v.size(), |
| 185 | + [] __device__(auto e) { return thrust::get<0>(e) == thrust::get<1>(e); })), |
| 186 | + handle.get_stream()); |
| 187 | + d_dst_v.resize(d_src_v.size(), handle.get_stream()); |
| 188 | + (*d_weight_v).resize(d_src_v.size(), handle.get_stream()); |
| 189 | + } else { |
| 190 | + auto edge_first = |
| 191 | + thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin())); |
| 192 | + d_src_v.resize( |
| 193 | + thrust::distance(edge_first, |
| 194 | + thrust::remove_if( |
| 195 | + handle.get_thrust_policy(), |
| 196 | + edge_first, |
| 197 | + edge_first + d_src_v.size(), |
| 198 | + [] __device__(auto e) { return thrust::get<0>(e) == thrust::get<1>(e); })), |
| 199 | + handle.get_stream()); |
| 200 | + d_dst_v.resize(d_src_v.size(), handle.get_stream()); |
| 201 | + } |
| 202 | + |
| 203 | + d_src_v.shrink_to_fit(handle.get_stream()); |
| 204 | + d_dst_v.shrink_to_fit(handle.get_stream()); |
| 205 | + if (d_weight_v) { (*d_weight_v).shrink_to_fit(handle.get_stream()); } |
| 206 | +} |
| 207 | + |
| 208 | +template void remove_self_loops( |
| 209 | + raft::handle_t const& handle, |
| 210 | + rmm::device_uvector<int32_t>& d_src_v /* [INOUT] */, |
| 211 | + rmm::device_uvector<int32_t>& d_dst_v /* [INOUT] */, |
| 212 | + std::optional<rmm::device_uvector<float>>& d_weight_v /* [INOUT] */); |
| 213 | + |
| 214 | +template void remove_self_loops( |
| 215 | + raft::handle_t const& handle, |
| 216 | + rmm::device_uvector<int32_t>& d_src_v /* [INOUT] */, |
| 217 | + rmm::device_uvector<int32_t>& d_dst_v /* [INOUT] */, |
| 218 | + std::optional<rmm::device_uvector<double>>& d_weight_v /* [INOUT] */); |
| 219 | + |
| 220 | +template void remove_self_loops( |
| 221 | + raft::handle_t const& handle, |
| 222 | + rmm::device_uvector<int64_t>& d_src_v /* [INOUT] */, |
| 223 | + rmm::device_uvector<int64_t>& d_dst_v /* [INOUT] */, |
| 224 | + std::optional<rmm::device_uvector<float>>& d_weight_v /* [INOUT] */); |
| 225 | + |
| 226 | +template void remove_self_loops( |
| 227 | + raft::handle_t const& handle, |
| 228 | + rmm::device_uvector<int64_t>& d_src_v /* [INOUT] */, |
| 229 | + rmm::device_uvector<int64_t>& d_dst_v /* [INOUT] */, |
| 230 | + std::optional<rmm::device_uvector<double>>& d_weight_v /* [INOUT] */); |
| 231 | + |
| 232 | +template <typename vertex_t, typename weight_t> |
| 233 | +void sort_and_remove_multi_edges( |
| 234 | + raft::handle_t const& handle, |
| 235 | + rmm::device_uvector<vertex_t>& d_src_v /* [INOUT] */, |
| 236 | + rmm::device_uvector<vertex_t>& d_dst_v /* [INOUT] */, |
| 237 | + std::optional<rmm::device_uvector<weight_t>>& d_weight_v /* [INOUT] */) |
| 238 | +{ |
| 239 | + if (d_weight_v) { |
| 240 | + auto edge_first = thrust::make_zip_iterator( |
| 241 | + thrust::make_tuple(d_src_v.begin(), d_dst_v.begin(), (*d_weight_v).begin())); |
| 242 | + thrust::sort(handle.get_thrust_policy(), edge_first, edge_first + d_src_v.size()); |
| 243 | + d_src_v.resize( |
| 244 | + thrust::distance(edge_first, |
| 245 | + thrust::unique(handle.get_thrust_policy(), |
| 246 | + edge_first, |
| 247 | + edge_first + d_src_v.size(), |
| 248 | + [] __device__(auto lhs, auto rhs) { |
| 249 | + return (thrust::get<0>(lhs) == thrust::get<0>(rhs)) && |
| 250 | + (thrust::get<1>(lhs) == thrust::get<1>(rhs)); |
| 251 | + })), |
| 252 | + handle.get_stream()); |
| 253 | + d_dst_v.resize(d_src_v.size(), handle.get_stream()); |
| 254 | + (*d_weight_v).resize(d_src_v.size(), handle.get_stream()); |
| 255 | + } else { |
| 256 | + auto edge_first = |
| 257 | + thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin())); |
| 258 | + thrust::sort(handle.get_thrust_policy(), edge_first, edge_first + d_src_v.size()); |
| 259 | + d_src_v.resize( |
| 260 | + thrust::distance( |
| 261 | + edge_first, |
| 262 | + thrust::unique(handle.get_thrust_policy(), edge_first, edge_first + d_src_v.size())), |
| 263 | + handle.get_stream()); |
| 264 | + d_dst_v.resize(d_src_v.size(), handle.get_stream()); |
| 265 | + } |
| 266 | + |
| 267 | + d_src_v.shrink_to_fit(handle.get_stream()); |
| 268 | + d_dst_v.shrink_to_fit(handle.get_stream()); |
| 269 | + if (d_weight_v) { (*d_weight_v).shrink_to_fit(handle.get_stream()); } |
| 270 | +} |
| 271 | + |
| 272 | +template void sort_and_remove_multi_edges( |
| 273 | + raft::handle_t const& handle, |
| 274 | + rmm::device_uvector<int32_t>& d_src_v /* [INOUT] */, |
| 275 | + rmm::device_uvector<int32_t>& d_dst_v /* [INOUT] */, |
| 276 | + std::optional<rmm::device_uvector<float>>& d_weight_v /* [INOUT] */); |
| 277 | + |
| 278 | +template void sort_and_remove_multi_edges( |
| 279 | + raft::handle_t const& handle, |
| 280 | + rmm::device_uvector<int32_t>& d_src_v /* [INOUT] */, |
| 281 | + rmm::device_uvector<int32_t>& d_dst_v /* [INOUT] */, |
| 282 | + std::optional<rmm::device_uvector<double>>& d_weight_v /* [INOUT] */); |
| 283 | + |
| 284 | +template void sort_and_remove_multi_edges( |
| 285 | + raft::handle_t const& handle, |
| 286 | + rmm::device_uvector<int64_t>& d_src_v /* [INOUT] */, |
| 287 | + rmm::device_uvector<int64_t>& d_dst_v /* [INOUT] */, |
| 288 | + std::optional<rmm::device_uvector<float>>& d_weight_v /* [INOUT] */); |
| 289 | + |
| 290 | +template void sort_and_remove_multi_edges( |
| 291 | + raft::handle_t const& handle, |
| 292 | + rmm::device_uvector<int64_t>& d_src_v /* [INOUT] */, |
| 293 | + rmm::device_uvector<int64_t>& d_dst_v /* [INOUT] */, |
| 294 | + std::optional<rmm::device_uvector<double>>& d_weight_v /* [INOUT] */); |
| 295 | + |
167 | 296 | } // namespace test |
168 | 297 | } // namespace cugraph |
0 commit comments