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.

Sign up for GitHub

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

Replace thrust::tuple with cuda::std::tuple #1546

Closed
wants to merge 2 commits into from
+239 −262
Closed

Replace thrust::tuple with cuda::std::tuple #1546

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
e8e0a93
Replace `thrust::tuple` with `cuda::std::tuple`
miscco Mar 4, 2025
7faefd0
Merge branch 'branch-25.04' into drop_thrust_tuple
bdice Mar 4, 2025
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
12 changes: 6 additions & 6 deletions cpp/doxygen/developer_guide/HEADER_ONLY_API_GUIDE.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -229,7 +229,7 @@ provided in `type_utils.hpp`.
So, to refactor the libcudf-based API, we remove the following code.

```c++
auto input_tuple = thrust::make_tuple(thrust::make_constant_iterator(static_cast<T>(radius)),
auto input_tuple = cuda::std::make_tuple(thrust::make_constant_iterator(static_cast<T>(radius)),
a_lon.begin<T>(),
a_lat.begin<T>(),
b_lon.begin<T>(),
Expand All @@ -242,11 +242,11 @@ thrust::transform(rmm::exec_policy(stream),
input_iter + result->size(),
result->mutable_view().begin<T>(),
[] __device__(auto inputs) {
return calculate_haversine_distance(thrust::get<0>(inputs),
thrust::get<1>(inputs),
thrust::get<2>(inputs),
thrust::get<3>(inputs),
thrust::get<4>(inputs));
return calculate_haversine_distance(cuda::std::get<0>(inputs),
cuda::std::get<1>(inputs),
cuda::std::get<2>(inputs),
cuda::std::get<3>(inputs),
cuda::std::get<4>(inputs));
});
```

Expand Down
14 changes: 7 additions & 7 deletions cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2024, NVIDIA CORPORATION.
* Copyright (c) 2022-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -84,13 +84,13 @@ CUSPATIAL_KERNEL void simple_find_and_combine_segments_kernel(OffsetRange offset
*/
template <typename index_t, typename T>
struct segment_comparator {
bool __device__ operator()(thrust::tuple<index_t, segment<T>> const& lhs,
thrust::tuple<index_t, segment<T>> const& rhs) const
bool __device__ operator()(cuda::std::tuple<index_t, segment<T>> const& lhs,
cuda::std::tuple<index_t, segment<T>> const& rhs) const
{
auto lhs_index = thrust::get<0>(lhs);
auto rhs_index = thrust::get<0>(rhs);
auto lhs_segment = thrust::get<1>(lhs);
auto rhs_segment = thrust::get<1>(rhs);
auto lhs_index = cuda::std::get<0>(lhs);
auto rhs_index = cuda::std::get<0>(rhs);
auto lhs_segment = cuda::std::get<1>(lhs);
auto rhs_segment = cuda::std::get<1>(rhs);

// Compare space id
if (lhs_index == rhs_index) {
Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cuspatial/detail/find/find_points_on_segments.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022, NVIDIA CORPORATION.
* Copyright (c) 2022-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand All @@ -26,8 +26,8 @@
#include <rmm/cuda_stream_view.hpp>
#include <rmm/exec_policy.hpp>

#include <cuda/std/tuple>
#include <thrust/tabulate.h>
#include <thrust/tuple.h>

namespace cuspatial {
namespace detail {
Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cuspatial/detail/functors.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2023, NVIDIA CORPORATION.
* Copyright (c) 2023-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -42,7 +42,7 @@ struct offset_pair_to_count_functor {
template <typename OffsetPairIterator>
CUSPATIAL_HOST_DEVICE auto operator()(OffsetPairIterator p)
{
return thrust::get<1>(p) - thrust::get<0>(p);
return cuda::std::get<1>(p) - cuda::std::get<0>(p);
}
};

Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cuspatial/detail/geometry/linestring_ref.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2024, NVIDIA CORPORATION.
* Copyright (c) 2022-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand All @@ -22,8 +22,8 @@
#include <cuspatial/iterator_factory.cuh>
#include <cuspatial/traits.hpp>

#include <cuda/std/tuple>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/tuple.h>

namespace cuspatial {

Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cuspatial/detail/geometry/polygon_ref.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2023-2024, NVIDIA CORPORATION.
* Copyright (c) 2023-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand All @@ -23,8 +23,8 @@
#include <cuspatial/iterator_factory.cuh>
#include <cuspatial/traits.hpp>

#include <cuda/std/tuple>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/tuple.h>

namespace cuspatial {

Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cuspatial/detail/index/construction/phase_1.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020-2024, NVIDIA CORPORATION.
* Copyright (c) 2020-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -27,6 +27,7 @@
#include <rmm/resource_ref.hpp>

#include <cuda/functional>
#include <cuda/std/tuple>
#include <thrust/copy.h>
#include <thrust/distance.h>
#include <thrust/fill.h>
Expand All @@ -38,7 +39,6 @@
#include <thrust/sequence.h>
#include <thrust/sort.h>
#include <thrust/transform.h>
#include <thrust/tuple.h>

#include <memory>
#include <tuple>
Expand Down
10 changes: 5 additions & 5 deletions cpp/include/cuspatial/detail/index/construction/phase_2.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020-2024, NVIDIA CORPORATION.
* Copyright (c) 2020-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand All @@ -23,6 +23,7 @@
#include <rmm/exec_policy.hpp>
#include <rmm/resource_ref.hpp>

#include <cuda/std/tuple>
#include <thrust/copy.h>
#include <thrust/count.h>
#include <thrust/distance.h>
Expand All @@ -38,7 +39,6 @@
#include <thrust/sequence.h>
#include <thrust/sort.h>
#include <thrust/transform.h>
#include <thrust/tuple.h>
#include <thrust/uninitialized_fill.h>

#include <memory>
Expand Down Expand Up @@ -85,9 +85,9 @@ inline rmm::device_uvector<uint32_t> flatten_point_keys(
keys_and_levels + num_valid_nodes,
flattened_keys.begin(),
[last_level = max_depth - 1] __device__(auto const& val) {
auto& key = thrust::get<0>(val);
auto& level = thrust::get<1>(val);
auto& is_parent = thrust::get<2>(val);
auto& key = cuda::std::get<0>(val);
auto& level = cuda::std::get<1>(val);
auto& is_parent = cuda::std::get<2>(val);
// if this is a parent node, return max_key. otherwise
// compute the key for one level up the tree. Leaf nodes
// whose keys are zero will be removed in a subsequent
Expand Down
12 changes: 6 additions & 6 deletions cpp/include/cuspatial/detail/index/construction/utilities.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020-2021, NVIDIA CORPORATION.
* Copyright (c) 2020-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand All @@ -18,18 +18,18 @@

#include <rmm/cuda_stream_view.hpp>

#include <thrust/tuple.h>
#include <cuda/std/tuple>

namespace cuspatial {
namespace detail {

template <typename T>
struct tuple_sum {
inline __device__ thrust::tuple<T, T> operator()(thrust::tuple<T, T> const& a,
thrust::tuple<T, T> const& b)
inline __device__ cuda::std::tuple<T, T> operator()(cuda::std::tuple<T, T> const& a,
cuda::std::tuple<T, T> const& b)
{
return thrust::make_tuple(thrust::get<0>(a) + thrust::get<0>(b),
thrust::get<1>(a) + thrust::get<1>(b));
return cuda::std::make_tuple(cuda::std::get<0>(a) + cuda::std::get<0>(b),
cuda::std::get<1>(a) + cuda::std::get<1>(b));
}
};

Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2024, NVIDIA CORPORATION.
* Copyright (c) 2022-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -37,6 +37,7 @@
#include <rmm/resource_ref.hpp>

#include <cuda/atomic>
#include <cuda/std/tuple>
#include <thrust/binary_search.h>
#include <thrust/iterator/constant_iterator.h>
#include <thrust/iterator/discard_iterator.h>
Expand All @@ -47,7 +48,6 @@
#include <thrust/sequence.h>
#include <thrust/sort.h>
#include <thrust/tabulate.h>
#include <thrust/tuple.h>
#include <thrust/uninitialized_fill.h>
#include <thrust/unique.h>

Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cuspatial/detail/intersection/linestring_intersection_count.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2024, NVIDIA CORPORATION.
* Copyright (c) 2022-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand All @@ -23,7 +23,7 @@

#include <rmm/cuda_stream_view.hpp>

#include <thrust/tuple.h>
#include <cuda/std/tuple>

#include <ranger/ranger.hpp>

Expand Down
8 changes: 4 additions & 4 deletions cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2024, NVIDIA CORPORATION.
* Copyright (c) 2022-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -31,14 +31,14 @@
#include <rmm/mr/device/device_memory_resource.hpp>
#include <rmm/resource_ref.hpp>

#include <cuda/std/tuple>
#include <thrust/binary_search.h>
#include <thrust/distance.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/reduce.h>
#include <thrust/remove.h>
#include <thrust/scan.h>
#include <thrust/tuple.h>
#include <thrust/uninitialized_fill.h>

#include <ranger/ranger.hpp>
Expand Down Expand Up @@ -162,9 +162,9 @@ struct id_ranges {

/// Row-wise getter to the id arrays
template <typename IndexType>
thrust::tuple<index_t, index_t, index_t, index_t> __device__ operator[](IndexType i)
cuda::std::tuple<index_t, index_t, index_t, index_t> __device__ operator[](IndexType i)
{
return thrust::make_tuple(
return cuda::std::make_tuple(
lhs_linestring_ids[i], lhs_segment_ids[i], rhs_linestring_ids[i], rhs_segment_ids[i]);
}

Expand Down
22 changes: 11 additions & 11 deletions cpp/include/cuspatial/detail/join/intersection.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020-2023, NVIDIA CORPORATION.
* Copyright (c) 2020-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand All @@ -24,14 +24,14 @@
#include <rmm/device_uvector.hpp>
#include <rmm/exec_policy.hpp>

#include <cuda/std/tuple>
#include <thrust/copy.h>
#include <thrust/distance.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/remove.h>
#include <thrust/transform.h>
#include <thrust/tuple.h>

#include <utility>

Expand All @@ -52,7 +52,7 @@ inline int32_t copy_leaf_intersections(InputIterator input_begin,
output_begin,
thrust::copy_if(
rmm::exec_policy(stream), input_begin, input_end, output_begin, [] __device__(auto const& t) {
return thrust::get<0>(t) == leaf_indicator;
return cuda::std::get<0>(t) == leaf_indicator;
}));
}

Expand All @@ -66,7 +66,7 @@ inline int32_t remove_non_quad_intersections(InputIterator input_begin,
output_begin,
thrust::remove_if(
rmm::exec_policy(stream), input_begin, input_end, output_begin, [] __device__(auto const& t) {
return thrust::get<0>(t) != quad_indicator;
return cuda::std::get<0>(t) != quad_indicator;
}));
}

Expand Down Expand Up @@ -98,13 +98,13 @@ inline std::pair<int32_t, int32_t> find_intersections(point_quadtree_ref quadtre
node_pairs,
[v_min, scale, max_depth, nodes = nodes_first, bboxes = bounding_box_first] __device__(
auto const& node_and_bbox) {
auto const& node_idx = thrust::get<0>(node_and_bbox);
auto const& bbox_idx = thrust::get<1>(node_and_bbox);
auto const& node_idx = cuda::std::get<0>(node_and_bbox);
auto const& bbox_idx = cuda::std::get<1>(node_and_bbox);

auto const& node = nodes[node_idx];
uint32_t const& key = thrust::get<0>(node);
uint8_t const& level = thrust::get<1>(node);
uint8_t const& is_internal_node = thrust::get<2>(node);
uint32_t const& key = cuda::std::get<0>(node);
uint8_t const& level = cuda::std::get<1>(node);
uint8_t const& is_internal_node = cuda::std::get<2>(node);

box<T> const bbox = bboxes[bbox_idx];
vec_2d<T> const bbox_min = bbox.v1;
Expand All @@ -121,10 +121,10 @@ inline std::pair<int32_t, int32_t> find_intersections(point_quadtree_ref quadtre
if ((node_x_min > bbox_max.x) || (node_x_max < bbox_min.x) || (node_y_min > bbox_max.y) ||
(node_y_max < bbox_min.y)) {
// if no overlap, return type = none_indicator
return thrust::make_tuple(none_indicator, level, node_idx, bbox_idx);
return cuda::std::make_tuple(none_indicator, level, node_idx, bbox_idx);
}
// otherwise return type = leaf_indicator (0) or quad_indicator (1)
return thrust::make_tuple(is_internal_node, level, node_idx, bbox_idx);
return cuda::std::make_tuple(is_internal_node, level, node_idx, bbox_idx);
});

auto num_leaves = copy_leaf_intersections(node_pairs, node_pairs + num_pairs, leaf_pairs, stream);
Expand Down
12 changes: 6 additions & 6 deletions cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2023-2024, NVIDIA CORPORATION.
* Copyright (c) 2023-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -58,14 +58,14 @@ struct compute_poly_and_point_indices {

using IndexType = iterator_value_type<QuadOffsetsIterator>;

inline thrust::tuple<IndexType, IndexType> __device__
inline cuda::std::tuple<IndexType, IndexType> __device__
operator()(std::uint64_t const global_index) const
{
auto const [quad_poly_index, local_point_index] =
get_quad_and_local_point_indices(global_index, point_offsets_begin, point_offsets_end);
auto const point_idx = quad_point_offsets_begin[quad_poly_index] + local_point_index;
auto const poly_idx = poly_indices_begin[quad_poly_index];
return thrust::make_tuple(poly_idx, point_idx);
return cuda::std::make_tuple(poly_idx, point_idx);
}
};

Expand All @@ -82,10 +82,10 @@ struct test_poly_point_intersection {
PointIterator points_first;
MultiPolygonRange polygons;

inline bool __device__ operator()(thrust::tuple<IndexType, IndexType> const& poly_point_idxs)
inline bool __device__ operator()(cuda::std::tuple<IndexType, IndexType> const& poly_point_idxs)
{
auto const poly_idx = thrust::get<0>(poly_point_idxs);
auto const point_idx = thrust::get<1>(poly_point_idxs);
auto const poly_idx = cuda::std::get<0>(poly_point_idxs);
auto const point_idx = cuda::std::get<1>(poly_point_idxs);

vec_2d<T> const& point = points_first[point_idx];

Expand Down
Loading
Loading