Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Adds load balanced kernel for point-point, point-linestring and linestring-linestring #1144

Draft
wants to merge 33 commits into
base: branch-23.08
Choose a base branch
from
Draft
Changes from 1 commit
Commits
Show all changes
33 commits
Select commit Hold shift + click to select a range
ac872e7
[skip ci] initial segment_counter working
isVoid May 11, 2023
84f317c
segment iterator works with empty linestrings
isVoid May 16, 2023
ce28b3c
enable polygon segment count test
isVoid May 16, 2023
cd2d229
add polygon segment methods
isVoid May 16, 2023
fcbe105
remove incomplete validations and add implicit assumptions in documen…
isVoid May 17, 2023
ee3576e
enable multipolygon range tests
isVoid May 17, 2023
7a53d05
improve doc of validators
isVoid May 18, 2023
2181e96
replace all multilinestring range tests with segment method test
isVoid May 18, 2023
5e3a45a
updates multipolygon tests to make use the new segment methods view
isVoid May 18, 2023
53e3ce0
remove segment methods from ranges and update all usage to `segment_m…
isVoid May 18, 2023
789f7a2
Cleanup: Remove all debug prints
isVoid May 18, 2023
2390acc
Cleanup: Remove segment functions in multilinestring_range and dead c…
isVoid May 18, 2023
728caa4
remove num_segments
isVoid May 18, 2023
5df7ea0
rename segment_methods -> multilinestring_segment
isVoid May 18, 2023
e030449
Updates segment_range and tests to make sure APIs exposed are all mul…
isVoid May 19, 2023
771f340
Refactors `linestring_polygon_distance` and add tests to include empt…
isVoid May 19, 2023
81e05bd
Move all segment range only functors to local file
isVoid May 19, 2023
9d00b41
Documentation for segment_range
isVoid May 19, 2023
f6255b8
typo
isVoid May 19, 2023
9ec9882
Merge branch 'branch-23.06' into fix/segment_iterator
isVoid May 19, 2023
0f151a5
Cleanups & doc improvement
isVoid May 19, 2023
57a2327
Merge branch 'fix/segment_iterator' of github.com:isVoid/cuspatial in…
isVoid May 19, 2023
a45f2c7
[skip ci] initial addition of the load balanced kernel
isVoid May 19, 2023
e6cd5d8
Merge branch 'fix/segment_iterator' into improvement/load_balanced_di…
isVoid May 22, 2023
8475fe8
Move detail iterator factory into detail folder
isVoid May 22, 2023
c463978
Remove functors.cuh, make iterator factory for it.
isVoid May 22, 2023
b442f7c
Remove every functors usage
isVoid May 22, 2023
8d76d0a
Impelement load balanced linestring distance
isVoid May 22, 2023
242987d
Merge branch 'branch-23.06' into improvement/load_balanced_distance_k…
isVoid May 26, 2023
64b638f
Merge branch 'branch-23.06' of https://github.com/rapidsai/cuspatial …
isVoid May 30, 2023
3621efd
optimize load_balanced_kernel
isVoid May 30, 2023
4633224
minor updates to balanced kernel
isVoid May 30, 2023
b8099d7
Merge branch 'improvement/load_balanced_distance_kernel' of github.co…
isVoid May 30, 2023
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
Cleanup: Remove segment functions in multilinestring_range and dead c…
…odes
isVoid committed May 18, 2023
commit 2390accb07921b06d5b8d16d772907ca0737e304
47 changes: 3 additions & 44 deletions cpp/include/cuspatial/detail/method/segment_method.cuh
Original file line number Diff line number Diff line change
@@ -39,16 +39,13 @@
namespace cuspatial {
namespace detail {

template <typename IndexType, IndexType value>
struct equals {
__device__ bool operator()(IndexType x) const { return x == value; }
};

template <typename IndexType>
struct greater_than_zero_functor {
__device__ IndexType operator()(IndexType x) const { return x > 0; }
};

// Optimization: for range that does not contain any empty linestrings,
// The _non_empty_linestring_prefix_sum can be initailized with counting_iterator.
template <typename ParentRange>
class segment_method {
using index_t = iterator_value_type<typename ParentRange::part_it_t>;
@@ -65,46 +62,11 @@ class segment_method {
thrust::make_zip_iterator(offset_range.begin(), thrust::next(offset_range.begin())),
offset_pair_to_count_functor{});

// // Preemptive test: does the given range contain any empty ring/linestring?
// _contains_empty_geom = thrust::any_of(
// rmm::exec_policy(stream),
// count_begin,
// count_begin + _range.num_linestrings(),
// equals<index_t, 0>{}
// );

// std::cout << std::boolalpha << "contains empty geometry: " << _contains_empty_geom <<
// std::endl;

auto count_greater_than_zero =
thrust::make_transform_iterator(count_begin, greater_than_zero_functor<index_t>{});

{
thrust::device_vector<index_t> count_greater_than_zero_truth(
count_greater_than_zero, count_greater_than_zero + _range.num_linestrings());

test::print_device_vector(count_greater_than_zero_truth, "count_greater_than_zero_truth: ");
}

// Compute the number of empty linestrings
// auto key_begin = make_geometry_id_iterator<index_t>(_range.geometry_offsets_begin(),
// _range.geometry_offsets_end());
// {
// thrust::device_vector<index_t> key_truth(key_begin, key_begin + _range.num_linestrings());

// test::print_device_vector(key_truth, "key_truth: ");
// }

zero_data_async(
_non_empty_linestring_prefix_sum.begin(), _non_empty_linestring_prefix_sum.end(), stream);
// thrust::reduce_by_key(rmm::exec_policy(stream),
// key_begin,
// key_begin + _range.num_linestrings(),
// count_greater_than_zero,
// thrust::make_discard_iterator(),
// thrust::next(_non_empty_linestring_prefix_sum.begin()),
// thrust::equal_to<index_t>{},
// thrust::plus<index_t>{});

thrust::inclusive_scan(rmm::exec_policy(stream),
count_greater_than_zero,
@@ -113,21 +75,18 @@ class segment_method {

_num_segments = _range.num_points() - _non_empty_linestring_prefix_sum.element(
_non_empty_linestring_prefix_sum.size() - 1, stream);

test::print_device_vector(_non_empty_linestring_prefix_sum, "non_empty_geometry_prefix_sum: ");
}

auto view()
{
auto index_range =
range(_non_empty_linestring_prefix_sum.begin(), _non_empty_linestring_prefix_sum.end());
return segment_method_view<ParentRange, decltype(index_range)>{
_range, index_range, _num_segments, _contains_empty_geom};
_range, index_range, _num_segments};
}

private:
ParentRange _range;
bool _contains_empty_geom;
index_t _num_segments;
rmm::device_uvector<index_t> _non_empty_linestring_prefix_sum;
};
26 changes: 3 additions & 23 deletions cpp/include/cuspatial/detail/method/segment_method_view.cuh
Original file line number Diff line number Diff line change
@@ -16,7 +16,6 @@

#pragma once

#include "thrust/iterator/permutation_iterator.h"
#include <cuspatial_test/test_util.cuh>

#include <cuspatial/cuda_utils.hpp>
@@ -28,6 +27,7 @@
#include <rmm/mr/device/per_device_resource.hpp>

#include <thrust/device_vector.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/iterator/transform_iterator.h>

namespace cuspatial {
@@ -40,32 +40,13 @@ class segment_method_view {
public:
segment_method_view(ParentRange range,
IndexRange non_empty_geometry_prefix_sum,
index_t num_segments,
bool contains_empty_ring)
index_t num_segments)
: _range(range),
_non_empty_geometry_prefix_sum(non_empty_geometry_prefix_sum),
_num_segments(num_segments),
_contains_empty_ring(contains_empty_ring)
_num_segments(num_segments)
{
thrust::device_vector<index_t> soffsets(segment_offset_begin(), segment_offset_end());
test::print_device_vector(soffsets, "Segment offsets: ");
}

// CUSPATIAL_HOST_DEVICE auto segment_count_begin()
// {
// auto num_points_begin = _range.multilinestring_point_count_begin();
// auto n_point_linestring_pair_it =
// thrust::make_zip_iterator(num_points_begin, this->_non_empty_linestring_count_begin());

// return thrust::make_transform_iterator(n_point_linestring_pair_it,
// point_count_to_segment_count_functor{});
// }

// CUSPATIAL_HOST_DEVICE auto segment_count_end()
// {
// return thrust::next(this->segment_count_begin(), _range.num_multilinestrings());
// }

CUSPATIAL_HOST_DEVICE index_t num_segments() { return _num_segments; }

CUSPATIAL_HOST_DEVICE auto segment_offset_begin()
@@ -113,7 +94,6 @@ class segment_method_view {
ParentRange _range;
IndexRange _non_empty_geometry_prefix_sum;
index_t _num_segments;
bool _contains_empty_ring;

CUSPATIAL_HOST_DEVICE auto _non_empty_linestring_count_begin()
{
38 changes: 3 additions & 35 deletions cpp/include/cuspatial/detail/range/multilinestring_range.cuh
Original file line number Diff line number Diff line change
@@ -25,13 +25,13 @@

#include <cuspatial/cuda_utils.hpp>
#include <cuspatial/detail/functors.cuh>
#include <cuspatial/detail/method/segment_method.cuh>
#include <cuspatial/detail/utility/validation.hpp>
#include <cuspatial/geometry/vec_2d.hpp>
#include <cuspatial/geometry_collection/multilinestring_ref.cuh>
#include <cuspatial/iterator_factory.cuh>
#include <cuspatial/range/multipoint_range.cuh>
#include <cuspatial/traits.hpp>
#include <cuspatial/detail/method/segment_method.cuh>

#include <thrust/iterator/permutation_iterator.h>

@@ -249,44 +249,12 @@ CUSPATIAL_HOST_DEVICE auto multilinestring_range<GeometryIterator, PartIterator,
}

template <typename GeometryIterator, typename PartIterator, typename VecIterator>
CUSPATIAL_HOST_DEVICE auto
multilinestring_range<GeometryIterator, PartIterator, VecIterator>::segment_begin()
{
return detail::make_counting_transform_iterator(
0,
detail::to_valid_segment_functor{
this->segment_offset_begin(), this->segment_offset_end(), thrust::make_counting_iterator(0), _point_begin});
}

template <typename GeometryIterator, typename PartIterator, typename VecIterator>
CUSPATIAL_HOST_DEVICE auto
multilinestring_range<GeometryIterator, PartIterator, VecIterator>::segment_end()
{
return segment_begin() + num_segments();
}

template <typename GeometryIterator, typename PartIterator, typename VecIterator>
auto
multilinestring_range<GeometryIterator, PartIterator, VecIterator>::segment_methods(rmm::cuda_stream_view stream)
auto multilinestring_range<GeometryIterator, PartIterator, VecIterator>::segment_methods(
rmm::cuda_stream_view stream)
{
return segment_method{*this, stream};
}

template <typename GeometryIterator, typename PartIterator, typename VecIterator>
CUSPATIAL_HOST_DEVICE auto
multilinestring_range<GeometryIterator, PartIterator, VecIterator>::segment_offset_begin()
{
return detail::make_counting_transform_iterator(0, detail::to_segment_offset_iterator{
_part_begin, thrust::make_counting_iterator(0)});
}

template <typename GeometryIterator, typename PartIterator, typename VecIterator>
CUSPATIAL_HOST_DEVICE auto
multilinestring_range<GeometryIterator, PartIterator, VecIterator>::segment_offset_end()
{
return segment_offset_begin() + thrust::distance(_part_begin, _part_end);
}

template <typename GeometryIterator, typename PartIterator, typename VecIterator>
CUSPATIAL_HOST_DEVICE auto
multilinestring_range<GeometryIterator, PartIterator, VecIterator>::as_multipoint_range()
9 changes: 0 additions & 9 deletions cpp/include/cuspatial/range/multilinestring_range.cuh
Original file line number Diff line number Diff line change
@@ -162,12 +162,6 @@ class multilinestring_range {
/// Returns an iterator to the counts of points per multilinestring
CUSPATIAL_HOST_DEVICE auto multilinestring_linestring_count_end();

/// Returns an iterator to the start of the segment
CUSPATIAL_HOST_DEVICE auto segment_begin();

/// Returns an iterator to the end of the segment
CUSPATIAL_HOST_DEVICE auto segment_end();

/// Constructs a segment methods object, can only be constructed on host.
/// To use segment methods on device, create a `segment_methods_view`
/// See: `segment_methods::view`
@@ -199,9 +193,6 @@ class multilinestring_range {
VecIterator _point_begin;
VecIterator _point_end;

CUSPATIAL_HOST_DEVICE auto segment_offset_begin();
CUSPATIAL_HOST_DEVICE auto segment_offset_end();

private:
/// @internal
/// Return the iterator to the part index where the point locates.