diff --git a/cpp/doxygen/developer_guide/HEADER_ONLY_API_GUIDE.md b/cpp/doxygen/developer_guide/HEADER_ONLY_API_GUIDE.md index 65bf1b7ab..f1bb86c40 100644 --- a/cpp/doxygen/developer_guide/HEADER_ONLY_API_GUIDE.md +++ b/cpp/doxygen/developer_guide/HEADER_ONLY_API_GUIDE.md @@ -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(radius)), +auto input_tuple = cuda::std::make_tuple(thrust::make_constant_iterator(static_cast(radius)), a_lon.begin(), a_lat.begin(), b_lon.begin(), @@ -242,11 +242,11 @@ thrust::transform(rmm::exec_policy(stream), input_iter + result->size(), result->mutable_view().begin(), [] __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)); }); ``` diff --git a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh index ca86096be..7bd5a9ca9 100644 --- a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh +++ b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh @@ -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. @@ -84,13 +84,13 @@ CUSPATIAL_KERNEL void simple_find_and_combine_segments_kernel(OffsetRange offset */ template struct segment_comparator { - bool __device__ operator()(thrust::tuple> const& lhs, - thrust::tuple> const& rhs) const + bool __device__ operator()(cuda::std::tuple> const& lhs, + cuda::std::tuple> 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) { diff --git a/cpp/include/cuspatial/detail/find/find_points_on_segments.cuh b/cpp/include/cuspatial/detail/find/find_points_on_segments.cuh index dfda0b8b1..d1887e624 100644 --- a/cpp/include/cuspatial/detail/find/find_points_on_segments.cuh +++ b/cpp/include/cuspatial/detail/find/find_points_on_segments.cuh @@ -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. @@ -26,8 +26,8 @@ #include #include +#include #include -#include namespace cuspatial { namespace detail { diff --git a/cpp/include/cuspatial/detail/functors.cuh b/cpp/include/cuspatial/detail/functors.cuh index 6690e6f24..736f26281 100644 --- a/cpp/include/cuspatial/detail/functors.cuh +++ b/cpp/include/cuspatial/detail/functors.cuh @@ -43,7 +43,7 @@ struct offset_pair_to_count_functor { template 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); } }; diff --git a/cpp/include/cuspatial/detail/geometry/linestring_ref.cuh b/cpp/include/cuspatial/detail/geometry/linestring_ref.cuh index 4b0d9d484..2cf8cceb7 100644 --- a/cpp/include/cuspatial/detail/geometry/linestring_ref.cuh +++ b/cpp/include/cuspatial/detail/geometry/linestring_ref.cuh @@ -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. @@ -22,8 +22,8 @@ #include #include +#include #include -#include namespace cuspatial { diff --git a/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh b/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh index a7da32644..1acdc4ff9 100644 --- a/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh +++ b/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh @@ -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. @@ -23,8 +23,8 @@ #include #include +#include #include -#include namespace cuspatial { diff --git a/cpp/include/cuspatial/detail/index/construction/phase_1.cuh b/cpp/include/cuspatial/detail/index/construction/phase_1.cuh index a444ae64a..4e6a40420 100644 --- a/cpp/include/cuspatial/detail/index/construction/phase_1.cuh +++ b/cpp/include/cuspatial/detail/index/construction/phase_1.cuh @@ -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. @@ -27,6 +27,7 @@ #include #include +#include #include #include #include @@ -38,7 +39,6 @@ #include #include #include -#include #include #include diff --git a/cpp/include/cuspatial/detail/index/construction/phase_2.cuh b/cpp/include/cuspatial/detail/index/construction/phase_2.cuh index df28ecca1..c9ba88394 100644 --- a/cpp/include/cuspatial/detail/index/construction/phase_2.cuh +++ b/cpp/include/cuspatial/detail/index/construction/phase_2.cuh @@ -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. @@ -23,6 +23,7 @@ #include #include +#include #include #include #include @@ -38,7 +39,6 @@ #include #include #include -#include #include #include @@ -85,9 +85,9 @@ inline rmm::device_uvector 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 diff --git a/cpp/include/cuspatial/detail/index/construction/utilities.cuh b/cpp/include/cuspatial/detail/index/construction/utilities.cuh index 987c46633..304078dce 100644 --- a/cpp/include/cuspatial/detail/index/construction/utilities.cuh +++ b/cpp/include/cuspatial/detail/index/construction/utilities.cuh @@ -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. @@ -18,18 +18,18 @@ #include -#include +#include namespace cuspatial { namespace detail { template struct tuple_sum { - inline __device__ thrust::tuple operator()(thrust::tuple const& a, - thrust::tuple const& b) + inline __device__ cuda::std::tuple operator()(cuda::std::tuple const& a, + cuda::std::tuple 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)); } }; diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh index 818c6560b..a8b81d408 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh @@ -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. @@ -37,6 +37,7 @@ #include #include +#include #include #include #include @@ -47,7 +48,6 @@ #include #include #include -#include #include #include diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection_count.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection_count.cuh index 19009c06c..8665269a5 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection_count.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection_count.cuh @@ -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. @@ -23,7 +23,7 @@ #include -#include +#include #include diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh index 5d2165437..8a04ccf98 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh @@ -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. @@ -31,6 +31,7 @@ #include #include +#include #include #include #include @@ -38,7 +39,6 @@ #include #include #include -#include #include #include @@ -162,9 +162,9 @@ struct id_ranges { /// Row-wise getter to the id arrays template - thrust::tuple __device__ operator[](IndexType i) + cuda::std::tuple __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]); } diff --git a/cpp/include/cuspatial/detail/join/intersection.cuh b/cpp/include/cuspatial/detail/join/intersection.cuh index cf137311e..fbb05ddec 100644 --- a/cpp/include/cuspatial/detail/join/intersection.cuh +++ b/cpp/include/cuspatial/detail/join/intersection.cuh @@ -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. @@ -24,6 +24,7 @@ #include #include +#include #include #include #include @@ -31,7 +32,6 @@ #include #include #include -#include #include @@ -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; })); } @@ -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; })); } @@ -98,13 +98,13 @@ inline std::pair 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 const bbox = bboxes[bbox_idx]; vec_2d const bbox_min = bbox.v1; @@ -121,10 +121,10 @@ inline std::pair 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); diff --git a/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh b/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh index 68a2e9596..4468ce665 100644 --- a/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh +++ b/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh @@ -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. @@ -58,14 +58,14 @@ struct compute_poly_and_point_indices { using IndexType = iterator_value_type; - inline thrust::tuple __device__ + inline cuda::std::tuple __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); } }; @@ -82,10 +82,10 @@ struct test_poly_point_intersection { PointIterator points_first; MultiPolygonRange polygons; - inline bool __device__ operator()(thrust::tuple const& poly_point_idxs) + inline bool __device__ operator()(cuda::std::tuple 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 const& point = points_first[point_idx]; diff --git a/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh b/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh index 5b61882b6..0af92ea81 100644 --- a/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh +++ b/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh @@ -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. @@ -124,7 +124,7 @@ struct compute_point_linestring_indices_and_distances { { } - inline __device__ thrust::tuple operator()(uint32_t const global_index) + inline __device__ cuda::std::tuple operator()(uint32_t const global_index) { auto const [point_id, linestring_id] = get_transposed_point_and_pair_index( global_index, point_offsets, point_offsets_end, quad_offsets, quad_offsets_end, quad_lengths); @@ -134,7 +134,7 @@ struct compute_point_linestring_indices_and_distances { auto const distance = point_linestring_distance(thrust::raw_reference_cast(points[point_id]), linestring); - return thrust::make_tuple(point_id, linestring_indices[linestring_id], distance); + return cuda::std::make_tuple(point_id, linestring_indices[linestring_id], distance); } }; @@ -247,10 +247,10 @@ quadtree_point_to_nearest_linestring(LinestringIndexIterator linestring_indices_ linestring_indices_first, linestrings}); - auto all_point_indices = - thrust::make_transform_iterator(all_point_linestring_indices_and_distances, - cuda::proclaim_return_type( - [] __device__(auto const& x) { return thrust::get<0>(x); })); + auto all_point_indices = thrust::make_transform_iterator( + all_point_linestring_indices_and_distances, + cuda::proclaim_return_type( + [] __device__(auto const& x) { return cuda::std::get<0>(x); })); // Allocate vectors for the distances min reduction auto num_points = std::distance(point_indices_first, point_indices_last); @@ -275,18 +275,18 @@ quadtree_point_to_nearest_linestring(LinestringIndexIterator linestring_indices_ thrust::make_discard_iterator(), output_linestring_idxs.begin(), output_distances.begin()), thrust::equal_to(), // comparator // binop to select the point/linestring pair with the smallest distance - cuda::proclaim_return_type>( + cuda::proclaim_return_type>( [] __device__(auto const& lhs, auto const& rhs) { - T const& d_lhs = thrust::get<2>(lhs); - T const& d_rhs = thrust::get<2>(rhs); + T const& d_lhs = cuda::std::get<2>(lhs); + T const& d_rhs = cuda::std::get<2>(rhs); // If lhs distance is 0, choose rhs if (d_lhs == T{0}) { return rhs; } // if rhs distance is 0, choose lhs if (d_rhs == T{0}) { return lhs; } // If distances to lhs/rhs are the same, choose linestring with smallest id if (d_lhs == d_rhs) { - auto const& i_lhs = thrust::get<1>(lhs); - auto const& i_rhs = thrust::get<1>(rhs); + auto const& i_lhs = cuda::std::get<1>(lhs); + auto const& i_rhs = cuda::std::get<1>(rhs); return i_lhs < i_rhs ? lhs : rhs; } // Otherwise choose linestring with smallest distance diff --git a/cpp/include/cuspatial/detail/join/traversal.cuh b/cpp/include/cuspatial/detail/join/traversal.cuh index 33cb2b98c..24f6a9b58 100644 --- a/cpp/include/cuspatial/detail/join/traversal.cuh +++ b/cpp/include/cuspatial/detail/join/traversal.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, 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. @@ -22,6 +22,7 @@ #include #include +#include #include #include #include @@ -32,7 +33,6 @@ #include #include #include -#include #include diff --git a/cpp/include/cuspatial/detail/pairwise_multipoint_equals_count.cuh b/cpp/include/cuspatial/detail/pairwise_multipoint_equals_count.cuh index 157c32485..a94a83c31 100644 --- a/cpp/include/cuspatial/detail/pairwise_multipoint_equals_count.cuh +++ b/cpp/include/cuspatial/detail/pairwise_multipoint_equals_count.cuh @@ -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. @@ -88,8 +88,7 @@ OutputIt pairwise_multipoint_equals_count(MultiPointRangeA lhs, thrust::copy( rmm::exec_policy(stream), rhs.point_begin(), rhs.point_end(), rhs_point_sorted.begin()); - auto rhs_with_keys = - thrust::make_zip_iterator(thrust::make_tuple(rhs_keys.begin(), rhs_point_sorted.begin())); + auto rhs_with_keys = thrust::make_zip_iterator(rhs_keys.begin(), rhs_point_sorted.begin()); thrust::sort(rmm::exec_policy(stream), rhs_with_keys, rhs_with_keys + rhs.num_points()); diff --git a/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh b/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh index 13270a7cc..d451d0a3f 100644 --- a/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh +++ b/cpp/include/cuspatial/detail/point_linestring_nearest_points.cuh @@ -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. @@ -99,19 +99,19 @@ CUSPATIAL_KERNEL void pairwise_point_linestring_nearest_points_kernel( auto distance_nearest_point_pair = point_to_segment_distance_squared_nearest_point(c, a, b); - auto distance_squared = thrust::get<0>(distance_nearest_point_pair); + auto distance_squared = cuda::std::get<0>(distance_nearest_point_pair); if (distance_squared < min_distance_squared) { min_distance_squared = distance_squared; nearest_point_idx = point_idx - point_start; nearest_part_idx = part_idx - linestring_parts_start; nearest_segment_idx = segment_idx - segment_start; - nearest_point = thrust::get<1>(distance_nearest_point_pair); + nearest_point = cuda::std::get<1>(distance_nearest_point_pair); } } } } - output_first[idx] = - thrust::make_tuple(nearest_point_idx, nearest_part_idx, nearest_segment_idx, nearest_point); + output_first[idx] = cuda::std::make_tuple( + nearest_point_idx, nearest_part_idx, nearest_segment_idx, nearest_point); } } diff --git a/cpp/include/cuspatial/detail/point_quadtree.cuh b/cpp/include/cuspatial/detail/point_quadtree.cuh index 19922af35..6ae8c6710 100644 --- a/cpp/include/cuspatial/detail/point_quadtree.cuh +++ b/cpp/include/cuspatial/detail/point_quadtree.cuh @@ -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. @@ -27,10 +27,10 @@ #include #include +#include #include #include #include -#include #include @@ -111,7 +111,7 @@ inline point_quadtree make_quad_tree(rmm::device_uvector& keys, offsets.begin(), // return is_internal_node ? lhs : rhs cuda::proclaim_return_type([] __device__(auto const& t) { - return thrust::get<0>(t) ? thrust::get<1>(t) : thrust::get<2>(t); + return cuda::std::get<0>(t) ? cuda::std::get<1>(t) : cuda::std::get<2>(t); })); return std::move(offsets); @@ -129,7 +129,7 @@ inline point_quadtree make_quad_tree(rmm::device_uvector& keys, lengths.begin(), // return bool ? lhs : rhs cuda::proclaim_return_type([] __device__(auto const& t) { - return thrust::get<0>(t) ? thrust::get<1>(t) : thrust::get<2>(t); + return cuda::std::get<0>(t) ? cuda::std::get<1>(t) : cuda::std::get<2>(t); })); // Shrink keys to the number of valid nodes diff --git a/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh b/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh index 11b8a594d..37edeb725 100644 --- a/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh +++ b/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh @@ -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. @@ -42,10 +42,10 @@ template struct trajectory_comparator { __device__ bool operator()(Tuple const& lhs, Tuple const& rhs) { - auto lhs_id = thrust::get<0>(lhs); - auto rhs_id = thrust::get<0>(rhs); - auto lhs_ts = thrust::get<1>(lhs); - auto rhs_ts = thrust::get<1>(rhs); + auto lhs_id = cuda::std::get<0>(lhs); + auto rhs_id = cuda::std::get<0>(rhs); + auto lhs_ts = cuda::std::get<1>(lhs); + auto rhs_ts = cuda::std::get<1>(rhs); return (lhs_id < rhs_id) || ((lhs_id == rhs_id) && (lhs_ts < rhs_ts)); }; }; @@ -68,7 +68,7 @@ void order_trajectories(IdInputIt ids_first, { using id_type = iterator_value_type; using timestamp_type = iterator_value_type; - using tuple_type = thrust::tuple; + using tuple_type = cuda::std::tuple; auto keys_first = thrust::make_zip_iterator(ids_first, timestamps_first); auto keys_out_first = thrust::make_zip_iterator(ids_out_first, timestamps_out_first); diff --git a/cpp/include/cuspatial/detail/trajectory/trajectory_distances_and_speeds.cuh b/cpp/include/cuspatial/detail/trajectory/trajectory_distances_and_speeds.cuh index 9ce7baef5..ae43e682e 100644 --- a/cpp/include/cuspatial/detail/trajectory/trajectory_distances_and_speeds.cuh +++ b/cpp/include/cuspatial/detail/trajectory/trajectory_distances_and_speeds.cuh @@ -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. @@ -80,16 +80,16 @@ OutputIt trajectory_distances_and_speeds(IndexT num_trajectories, id_point_timestamp + 1, duration_and_distance + 1, [] __device__(auto const& p0, auto const& p1) { - if (thrust::get<0>(p0) == thrust::get<0>(p1)) { // ids are the same - Point pos0 = thrust::get<1>(p0); - Point pos1 = thrust::get<1>(p1); - Timestamp t0 = thrust::get<2>(p0); - Timestamp t1 = thrust::get<2>(p1); + if (cuda::std::get<0>(p0) == cuda::std::get<0>(p1)) { // ids are the same + Point pos0 = cuda::std::get<1>(p0); + Point pos1 = cuda::std::get<1>(p1); + Timestamp t0 = cuda::std::get<2>(p0); + Timestamp t1 = cuda::std::get<2>(p1); Point vec = pos1 - pos0; // duration and distance - return thrust::make_tuple((t1 - t0).count(), sqrt(dot(vec, vec))); + return cuda::std::make_tuple((t1 - t0).count(), sqrt(dot(vec, vec))); } - return thrust::make_tuple(Rep{}, T{}); + return cuda::std::make_tuple(Rep{}, T{}); }); auto duration_and_distance_tmp = thrust::make_zip_iterator(durations.begin(), @@ -100,8 +100,8 @@ OutputIt trajectory_distances_and_speeds(IndexT num_trajectories, rmm::device_uvector durations_tmp(num_trajectories, stream); rmm::device_uvector distances_tmp(num_trajectories, stream); - auto distances_begin = thrust::get<0>(distances_and_speeds_first.get_iterator_tuple()); - auto speeds_begin = thrust::get<1>(distances_and_speeds_first.get_iterator_tuple()); + auto distances_begin = cuda::std::get<0>(distances_and_speeds_first.get_iterator_tuple()); + auto speeds_begin = cuda::std::get<1>(distances_and_speeds_first.get_iterator_tuple()); auto duration_distances_and_speed = thrust::make_zip_iterator( durations_tmp.begin(), distances_tmp.begin(), distances_begin, speeds_begin); @@ -119,13 +119,13 @@ OutputIt trajectory_distances_and_speeds(IndexT num_trajectories, duration_distances_and_speed, thrust::equal_to(), [] __device__(auto a, auto b) { - auto time_d = Dur(thrust::get<0>(a)) + Dur(thrust::get<0>(b)); + auto time_d = Dur(cuda::std::get<0>(a)) + Dur(cuda::std::get<0>(b)); auto time_s = static_cast(time_d.count()) * static_cast(Period::num) / static_cast(Period::den); - T dist_km = thrust::get<1>(a) + thrust::get<1>(b); + T dist_km = cuda::std::get<1>(a) + cuda::std::get<1>(b); T dist_m = dist_km * T{1000.0}; // km to m T speed_m_s = dist_m / time_s; // m/ms to m/s - return thrust::make_tuple(time_d.count(), dist_km, dist_m, speed_m_s); + return cuda::std::make_tuple(time_d.count(), dist_km, dist_m, speed_m_s); }); // check for errors diff --git a/cpp/include/cuspatial/detail/utility/linestring.cuh b/cpp/include/cuspatial/detail/utility/linestring.cuh index 261648ac1..c055592d1 100644 --- a/cpp/include/cuspatial/detail/utility/linestring.cuh +++ b/cpp/include/cuspatial/detail/utility/linestring.cuh @@ -21,9 +21,9 @@ #include #include +#include #include #include -#include namespace cuspatial { namespace detail { @@ -52,7 +52,7 @@ endpoint_index_of_linestring(SizeType const& linestring_idx, * ab */ template -__forceinline__ thrust::tuple> __device__ +__forceinline__ cuda::std::tuple> __device__ point_to_segment_distance_squared_nearest_point(vec_2d const& c, vec_2d const& a, vec_2d const& b) @@ -60,18 +60,18 @@ point_to_segment_distance_squared_nearest_point(vec_2d const& c, auto ab = b - a; auto ac = c - a; auto l_squared = dot(ab, ab); - if (float_equal(l_squared, T{0})) { return thrust::make_tuple(dot(ac, ac), a); } + if (float_equal(l_squared, T{0})) { return cuda::std::make_tuple(dot(ac, ac), a); } auto r = dot(ac, ab); auto bc = c - b; // If the projection of `c` is outside of segment `ab`, compute point-point distance. if (r <= 0 or r >= l_squared) { auto dac = dot(ac, ac); auto dbc = dot(bc, bc); - return dac < dbc ? thrust::make_tuple(dac, a) : thrust::make_tuple(dbc, b); + return dac < dbc ? cuda::std::make_tuple(dac, a) : cuda::std::make_tuple(dbc, b); } auto p = a + (r / l_squared) * ab; auto pc = c - p; - return thrust::make_tuple(dot(pc, pc), p); + return cuda::std::make_tuple(dot(pc, pc), p); } /** diff --git a/cpp/include/cuspatial/iterator_factory.cuh b/cpp/include/cuspatial/iterator_factory.cuh index 9b22ba9bc..3093dea1b 100644 --- a/cpp/include/cuspatial/iterator_factory.cuh +++ b/cpp/include/cuspatial/iterator_factory.cuh @@ -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. @@ -22,6 +22,7 @@ #include #include +#include #include #include #include @@ -29,7 +30,6 @@ #include #include #include -#include #include #include @@ -53,9 +53,9 @@ inline CUSPATIAL_HOST_DEVICE auto make_counting_transform_iterator(IndexType sta */ template > struct tuple_to_vec_2d { - __device__ VectorType operator()(thrust::tuple const& pos) + __device__ VectorType operator()(cuda::std::tuple const& pos) { - return VectorType{thrust::get<0>(pos), thrust::get<1>(pos)}; + return VectorType{cuda::std::get<0>(pos), cuda::std::get<1>(pos)}; } }; @@ -65,9 +65,9 @@ struct tuple_to_vec_2d { */ template > struct vec_2d_to_tuple { - __device__ thrust::tuple operator()(VectorType const& xy) + __device__ cuda::std::tuple operator()(VectorType const& xy) { - return thrust::make_tuple(xy.x, xy.y); + return cuda::std::make_tuple(xy.x, xy.y); } }; @@ -77,9 +77,9 @@ struct vec_2d_to_tuple { */ template > struct box_to_tuple { - __device__ thrust::tuple operator()(Box const& box) + __device__ cuda::std::tuple operator()(Box const& box) { - return thrust::make_tuple(box.v1.x, box.v1.y, box.v2.x, box.v2.y); + return cuda::std::make_tuple(box.v1.x, box.v1.y, box.v2.x, box.v2.y); } }; @@ -89,10 +89,10 @@ struct box_to_tuple { */ template > struct vec_2d_tuple_to_box { - __device__ box operator()(thrust::tuple pair) + __device__ box operator()(cuda::std::tuple pair) { - auto v1 = thrust::get<0>(pair); - auto v2 = thrust::get<1>(pair); + auto v1 = cuda::std::get<0>(pair); + auto v2 = cuda::std::get<1>(pair); return {v1, v2}; } }; @@ -258,7 +258,7 @@ template auto make_vec_2d_output_iterator(FirstIter first, SecondIter second) { using T = typename std::iterator_traits::value_type; - auto zipped_out = thrust::make_zip_iterator(thrust::make_tuple(first, second)); + auto zipped_out = thrust::make_zip_iterator(first, second); return thrust::transform_output_iterator(zipped_out, detail::vec_2d_to_tuple()); } @@ -280,8 +280,7 @@ auto make_vec_2d_output_iterator(Iter d_points_begin) auto odd_positions = thrust::make_permutation_iterator( thrust::next(d_points_begin), detail::make_counting_transform_iterator(0, fixed_stride_2_functor)); - auto zipped_outputs = - thrust::make_zip_iterator(thrust::make_tuple(even_positions, odd_positions)); + auto zipped_outputs = thrust::make_zip_iterator(even_positions, odd_positions); return thrust::make_transform_output_iterator(zipped_outputs, detail::vec_2d_to_tuple()); } diff --git a/cpp/include/cuspatial/point_quadtree.cuh b/cpp/include/cuspatial/point_quadtree.cuh index f64cee6f0..c5efb97ce 100644 --- a/cpp/include/cuspatial/point_quadtree.cuh +++ b/cpp/include/cuspatial/point_quadtree.cuh @@ -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. @@ -24,8 +24,8 @@ #include #include +#include #include -#include #include diff --git a/cpp/include/cuspatial_test/vector_equality.hpp b/cpp/include/cuspatial_test/vector_equality.hpp index a8278f2d0..f5fbf9adc 100644 --- a/cpp/include/cuspatial_test/vector_equality.hpp +++ b/cpp/include/cuspatial_test/vector_equality.hpp @@ -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. @@ -224,7 +224,7 @@ std::pair>, rmm::device_vector>> unpack_v thrust::transform(pairs.begin(), pairs.end(), zipped_output, [] __device__(Pair const& pair) { auto [a, b] = pair; - return thrust::make_tuple(a, b); + return cuda::std::make_tuple(a, b); }); return {std::move(first), std::move(second)}; } diff --git a/cpp/src/equality/pairwise_multipoint_equals_count.cu b/cpp/src/equality/pairwise_multipoint_equals_count.cu index abfde9e9c..84e006bed 100644 --- a/cpp/src/equality/pairwise_multipoint_equals_count.cu +++ b/cpp/src/equality/pairwise_multipoint_equals_count.cu @@ -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. @@ -28,10 +28,10 @@ #include #include +#include #include #include #include -#include #include #include diff --git a/cpp/src/nearest_points/point_linestring_nearest_points.cu b/cpp/src/nearest_points/point_linestring_nearest_points.cu index 00a4d365f..5142e1a87 100644 --- a/cpp/src/nearest_points/point_linestring_nearest_points.cu +++ b/cpp/src/nearest_points/point_linestring_nearest_points.cu @@ -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. @@ -87,10 +87,10 @@ struct pairwise_point_linestring_nearest_points_impl { if constexpr (!is_multi_point && !is_multi_linestring) { auto output_its = thrust::make_zip_iterator( - thrust::make_tuple(thrust::make_discard_iterator(), - thrust::make_discard_iterator(), - segment_idx->mutable_view().begin(), - nearest_points_it)); + cuda::std::make_tuple(thrust::make_discard_iterator(), + thrust::make_discard_iterator(), + segment_idx->mutable_view().begin(), + nearest_points_it)); pairwise_point_linestring_nearest_points(point_geometry_it, point_geometry_it + num_pairs + 1, @@ -114,10 +114,10 @@ struct pairwise_point_linestring_nearest_points_impl { stream, mr); auto output_its = thrust::make_zip_iterator( - thrust::make_tuple(nearest_point_idx->mutable_view().begin(), - thrust::make_discard_iterator(), - segment_idx->mutable_view().begin(), - nearest_points_it)); + cuda::std::make_tuple(nearest_point_idx->mutable_view().begin(), + thrust::make_discard_iterator(), + segment_idx->mutable_view().begin(), + nearest_points_it)); pairwise_point_linestring_nearest_points(point_geometry_it, point_geometry_it + num_pairs + 1, @@ -143,10 +143,10 @@ struct pairwise_point_linestring_nearest_points_impl { stream, mr); auto output_its = thrust::make_zip_iterator( - thrust::make_tuple(thrust::make_discard_iterator(), - nearest_linestring_idx->mutable_view().begin(), - segment_idx->mutable_view().begin(), - nearest_points_it)); + cuda::std::make_tuple(thrust::make_discard_iterator(), + nearest_linestring_idx->mutable_view().begin(), + segment_idx->mutable_view().begin(), + nearest_points_it)); pairwise_point_linestring_nearest_points(point_geometry_it, point_geometry_it + num_pairs + 1, @@ -178,10 +178,10 @@ struct pairwise_point_linestring_nearest_points_impl { stream, mr); auto output_its = thrust::make_zip_iterator( - thrust::make_tuple(nearest_point_idx->mutable_view().begin(), - nearest_linestring_idx->mutable_view().begin(), - segment_idx->mutable_view().begin(), - nearest_points_it)); + cuda::std::make_tuple(nearest_point_idx->mutable_view().begin(), + nearest_linestring_idx->mutable_view().begin(), + segment_idx->mutable_view().begin(), + nearest_points_it)); pairwise_point_linestring_nearest_points(point_geometry_it, point_geometry_it + num_pairs + 1, diff --git a/cpp/src/projection/sinusoidal_projection.cu b/cpp/src/projection/sinusoidal_projection.cu index 2552601a4..0483de73d 100644 --- a/cpp/src/projection/sinusoidal_projection.cu +++ b/cpp/src/projection/sinusoidal_projection.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2024, NVIDIA CORPORATION. + * Copyright (c) 2019-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. @@ -28,10 +28,10 @@ #include #include +#include #include #include #include -#include #include #include diff --git a/cpp/src/trajectory/trajectory_distances_and_speeds.cu b/cpp/src/trajectory/trajectory_distances_and_speeds.cu index 89474e925..858e93562 100644 --- a/cpp/src/trajectory/trajectory_distances_and_speeds.cu +++ b/cpp/src/trajectory/trajectory_distances_and_speeds.cu @@ -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. @@ -31,6 +31,7 @@ #include #include +#include #include #include #include @@ -39,7 +40,6 @@ #include #include #include -#include namespace cuspatial { diff --git a/cpp/tests/distance/point_distance_test.cu b/cpp/tests/distance/point_distance_test.cu index 72476ee19..0654f04c6 100644 --- a/cpp/tests/distance/point_distance_test.cu +++ b/cpp/tests/distance/point_distance_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, 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. @@ -28,6 +28,7 @@ #include #include +#include #include #include #include @@ -36,7 +37,6 @@ #include #include #include -#include #include #include @@ -99,11 +99,10 @@ auto compute_point_distance_host(Cart2DVec const& point1, Cart2DVec const& point using T = typename Cart2D::value_type; thrust::host_vector h_point1(point1); thrust::host_vector h_point2(point2); - auto pair_iter = - thrust::make_zip_iterator(thrust::make_tuple(h_point1.begin(), h_point2.begin())); + auto pair_iter = thrust::make_zip_iterator(h_point1.begin(), h_point2.begin()); auto result_iter = thrust::make_transform_iterator(pair_iter, [](auto p) { - auto p0 = thrust::get<0>(p); - auto p1 = thrust::get<1>(p); + auto p0 = cuda::std::get<0>(p); + auto p1 = cuda::std::get<1>(p); return std::sqrt(dot(p0 - p1, p0 - p1)); }); diff --git a/cpp/tests/intersection/intersection_test_utils.cuh b/cpp/tests/intersection/intersection_test_utils.cuh index 3a2112d07..c355a5c86 100644 --- a/cpp/tests/intersection/intersection_test_utils.cuh +++ b/cpp/tests/intersection/intersection_test_utils.cuh @@ -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. @@ -24,11 +24,11 @@ #include +#include #include #include #include #include -#include namespace cuspatial { namespace test { @@ -55,13 +55,13 @@ bool CUSPATIAL_HOST_DEVICE operator<(segment lhs, segment rhs) */ template struct order_key_value_pairs { - using key_value_t = thrust::tuple; + using key_value_t = cuda::std::tuple; bool CUSPATIAL_HOST_DEVICE operator()(key_value_t lhs, key_value_t rhs) { - return thrust::get<0>(lhs) < thrust::get<0>(rhs) || - (thrust::get<0>(lhs) == thrust::get<0>(rhs) && - thrust::get<1>(lhs) < thrust::get<1>(rhs)); + return cuda::std::get<0>(lhs) < cuda::std::get<0>(rhs) || + (cuda::std::get<0>(lhs) == cuda::std::get<0>(rhs) && + cuda::std::get<1>(lhs) < cuda::std::get<1>(rhs)); } }; @@ -133,7 +133,7 @@ linestring_intersection_result segment_sort_intersection_result( keys_points_begin, keys_points_begin + num_points, scatter_map.begin(), - order_key_value_pairs, vec_2d>{}); + order_key_value_pairs, vec_2d>{}); // Segment-sort the segment array auto keys_segment_begin = @@ -143,7 +143,7 @@ linestring_intersection_result segment_sort_intersection_result( keys_segment_begin, keys_segment_begin + num_segments, scatter_map.begin() + num_points, - order_key_value_pairs, segment>{}); + order_key_value_pairs, segment>{}); // Restore the order of indices auto lhs_linestring_id = std::make_unique>(num_geoms, stream, mr); diff --git a/cpp/tests/nearest_points/point_linestring_nearest_points_test.cu b/cpp/tests/nearest_points/point_linestring_nearest_points_test.cu index dfeac5cb2..a6301ef08 100644 --- a/cpp/tests/nearest_points/point_linestring_nearest_points_test.cu +++ b/cpp/tests/nearest_points/point_linestring_nearest_points_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, 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. @@ -58,11 +58,10 @@ TYPED_TEST(PairwisePointLinestringNearestPointsTest, Empty) auto nearest_linestring_segment_id = rmm::device_vector(0); auto neartest_point_coordinate = rmm::device_vector>(0); - auto output_it = - thrust::make_zip_iterator(thrust::make_tuple(nearest_point_id.begin(), - nearest_linestring_parts_id.begin(), - nearest_linestring_segment_id.begin(), - neartest_point_coordinate.begin())); + auto output_it = thrust::make_zip_iterator(nearest_point_id.begin(), + nearest_linestring_parts_id.begin(), + nearest_linestring_segment_id.begin(), + neartest_point_coordinate.begin()); auto ret = pairwise_point_linestring_nearest_points(points_geometry_it, points_geometry_it + num_pairs + 1, @@ -98,11 +97,10 @@ TYPED_TEST(PairwisePointLinestringNearestPointsTest, OnePairSingleComponent) auto nearest_linestring_segment_id = rmm::device_vector(num_pairs); auto neartest_point_coordinate = rmm::device_vector>(num_pairs); - auto output_it = - thrust::make_zip_iterator(thrust::make_tuple(thrust::make_discard_iterator(), - thrust::make_discard_iterator(), - nearest_linestring_segment_id.begin(), - neartest_point_coordinate.begin())); + auto output_it = thrust::make_zip_iterator(thrust::make_discard_iterator(), + thrust::make_discard_iterator(), + nearest_linestring_segment_id.begin(), + neartest_point_coordinate.begin()); auto ret = pairwise_point_linestring_nearest_points(points_geometry_it, points_geometry_it + num_pairs + 1, @@ -137,11 +135,10 @@ TYPED_TEST(PairwisePointLinestringNearestPointsTest, NearestAtLeftEndPoint) auto nearest_linestring_segment_id = rmm::device_vector(num_pairs); auto neartest_point_coordinate = rmm::device_vector>(num_pairs); - auto output_it = - thrust::make_zip_iterator(thrust::make_tuple(thrust::make_discard_iterator(), - thrust::make_discard_iterator(), - nearest_linestring_segment_id.begin(), - neartest_point_coordinate.begin())); + auto output_it = thrust::make_zip_iterator(thrust::make_discard_iterator(), + thrust::make_discard_iterator(), + nearest_linestring_segment_id.begin(), + neartest_point_coordinate.begin()); auto ret = pairwise_point_linestring_nearest_points(points_geometry_it, points_geometry_it + num_pairs + 1, @@ -176,11 +173,10 @@ TYPED_TEST(PairwisePointLinestringNearestPointsTest, NearestAtRightEndPoint) auto nearest_linestring_segment_id = rmm::device_vector(num_pairs); auto neartest_point_coordinate = rmm::device_vector>(num_pairs); - auto output_it = - thrust::make_zip_iterator(thrust::make_tuple(thrust::make_discard_iterator(), - thrust::make_discard_iterator(), - nearest_linestring_segment_id.begin(), - neartest_point_coordinate.begin())); + auto output_it = thrust::make_zip_iterator(thrust::make_discard_iterator(), + thrust::make_discard_iterator(), + nearest_linestring_segment_id.begin(), + neartest_point_coordinate.begin()); auto ret = pairwise_point_linestring_nearest_points(points_geometry_it, points_geometry_it + num_pairs + 1, @@ -216,11 +212,10 @@ TYPED_TEST(PairwisePointLinestringNearestPointsTest, PointAtEndPoints) auto nearest_linestring_segment_id = rmm::device_vector(num_pairs); auto neartest_point_coordinate = rmm::device_vector>(num_pairs); - auto output_it = - thrust::make_zip_iterator(thrust::make_tuple(thrust::make_discard_iterator(), - thrust::make_discard_iterator(), - nearest_linestring_segment_id.begin(), - neartest_point_coordinate.begin())); + auto output_it = thrust::make_zip_iterator(thrust::make_discard_iterator(), + thrust::make_discard_iterator(), + nearest_linestring_segment_id.begin(), + neartest_point_coordinate.begin()); auto ret = pairwise_point_linestring_nearest_points(points_geometry_it, points_geometry_it + num_pairs + 1, @@ -256,11 +251,10 @@ TYPED_TEST(PairwisePointLinestringNearestPointsTest, PointOnLineString) auto nearest_linestring_segment_id = rmm::device_vector(num_pairs); auto neartest_point_coordinate = rmm::device_vector>(num_pairs); - auto output_it = - thrust::make_zip_iterator(thrust::make_tuple(thrust::make_discard_iterator(), - thrust::make_discard_iterator(), - nearest_linestring_segment_id.begin(), - neartest_point_coordinate.begin())); + auto output_it = thrust::make_zip_iterator(thrust::make_discard_iterator(), + thrust::make_discard_iterator(), + nearest_linestring_segment_id.begin(), + neartest_point_coordinate.begin()); auto ret = pairwise_point_linestring_nearest_points(points_geometry_it, points_geometry_it + num_pairs + 1, @@ -296,11 +290,10 @@ TYPED_TEST(PairwisePointLinestringNearestPointsTest, TwoPairsSingleComponent) auto nearest_linestring_segment_id = rmm::device_vector(num_pairs); auto neartest_point_coordinate = rmm::device_vector>(num_pairs); - auto output_it = - thrust::make_zip_iterator(thrust::make_tuple(thrust::make_discard_iterator(), - thrust::make_discard_iterator(), - nearest_linestring_segment_id.begin(), - neartest_point_coordinate.begin())); + auto output_it = thrust::make_zip_iterator(thrust::make_discard_iterator(), + thrust::make_discard_iterator(), + nearest_linestring_segment_id.begin(), + neartest_point_coordinate.begin()); auto ret = pairwise_point_linestring_nearest_points(points_geometry_it, points_geometry_it + num_pairs + 1, @@ -338,11 +331,10 @@ TYPED_TEST(PairwisePointLinestringNearestPointsTest, OnePairMultiComponent) auto nearest_linestring_segment_id = rmm::device_vector(num_pairs); auto neartest_point_coordinate = rmm::device_vector>(num_pairs); - auto output_it = - thrust::make_zip_iterator(thrust::make_tuple(nearest_point_id.begin(), - nearest_linestring_linestring_id.begin(), - nearest_linestring_segment_id.begin(), - neartest_point_coordinate.begin())); + auto output_it = thrust::make_zip_iterator(nearest_point_id.begin(), + nearest_linestring_linestring_id.begin(), + nearest_linestring_segment_id.begin(), + neartest_point_coordinate.begin()); auto ret = pairwise_point_linestring_nearest_points(points_geometry_offsets.begin(), points_geometry_offsets.end(), @@ -395,10 +387,10 @@ TYPED_TEST(PairwisePointLinestringNearestPointsTest, ThreePairMultiComponent) auto nearest_segment_id = rmm::device_vector(num_pairs); auto neartest_point_coordinate = rmm::device_vector>(num_pairs); - auto output_it = thrust::make_zip_iterator(thrust::make_tuple(nearest_point_id.begin(), - nearest_linestring_id.begin(), - nearest_segment_id.begin(), - neartest_point_coordinate.begin())); + auto output_it = thrust::make_zip_iterator(nearest_point_id.begin(), + nearest_linestring_id.begin(), + nearest_segment_id.begin(), + neartest_point_coordinate.begin()); auto ret = pairwise_point_linestring_nearest_points(points_geometry_offsets.begin(), points_geometry_offsets.end(), diff --git a/cpp/tests/operators/linestrings_test.cu b/cpp/tests/operators/linestrings_test.cu index 66c4d0ec6..6950cff4a 100644 --- a/cpp/tests/operators/linestrings_test.cu +++ b/cpp/tests/operators/linestrings_test.cu @@ -81,14 +81,14 @@ CUSPATIAL_KERNEL void compute_intersection(segment ab, template struct unpack_optional_segment { - thrust::tuple, optional_vec2d> CUSPATIAL_HOST_DEVICE + cuda::std::tuple, optional_vec2d> CUSPATIAL_HOST_DEVICE operator()(cuda::std::optional> segment) { if (segment.has_value()) - return thrust::make_tuple(segment.value().v1, segment.value().v2); + return cuda::std::make_tuple(segment.value().v1, segment.value().v2); else - return thrust::tuple, optional_vec2d>{cuda::std::nullopt, - cuda::std::nullopt}; + return cuda::std::tuple, optional_vec2d>{cuda::std::nullopt, + cuda::std::nullopt}; } }; diff --git a/cpp/tests/trajectory/trajectory_test_utils.cuh b/cpp/tests/trajectory/trajectory_test_utils.cuh index b150e7284..98630edce 100644 --- a/cpp/tests/trajectory/trajectory_test_utils.cuh +++ b/cpp/tests/trajectory/trajectory_test_utils.cuh @@ -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. @@ -227,15 +227,15 @@ struct trajectory_test_data { } struct duration_functor { - using id_and_timestamp = thrust::tuple; + using id_and_timestamp = cuda::std::tuple; __host__ __device__ time_point::rep operator()(id_and_timestamp const& p0, id_and_timestamp const& p1) { - auto const id0 = thrust::get<0>(p0); - auto const id1 = thrust::get<0>(p1); - auto const t0 = thrust::get<1>(p0); - auto const t1 = thrust::get<1>(p1); + auto const id0 = cuda::std::get<0>(p0); + auto const id1 = cuda::std::get<0>(p1); + auto const t0 = cuda::std::get<1>(p0); + auto const t1 = cuda::std::get<1>(p1); if (id0 == id1) { return (t1 - t0).count(); } return 0; @@ -243,14 +243,14 @@ struct trajectory_test_data { }; struct distance_functor { - using id_and_position = thrust::tuple>; + using id_and_position = cuda::std::tuple>; __host__ __device__ T operator()(id_and_position const& p0, id_and_position const& p1) { - auto const id0 = thrust::get<0>(p0); - auto const id1 = thrust::get<0>(p1); + auto const id0 = cuda::std::get<0>(p0); + auto const id1 = cuda::std::get<0>(p1); if (id0 == id1) { - auto const pos0 = thrust::get<1>(p0); - auto const pos1 = thrust::get<1>(p1); + auto const pos0 = cuda::std::get<1>(p0); + auto const pos1 = cuda::std::get<1>(p1); auto const vec = pos1 - pos0; return sqrt(dot(vec, vec)); } @@ -259,7 +259,7 @@ struct trajectory_test_data { }; struct average_distance_speed_functor { - using duration_distance = thrust::tuple; + using duration_distance = cuda::std::tuple; using Sec = typename cuda::std::chrono::seconds; using Period = typename cuda::std::ratio_divide::type; @@ -268,10 +268,10 @@ struct trajectory_test_data { duration_distance const& b) { auto time_d = - time_point::duration(thrust::get<0>(a)) + time_point::duration(thrust::get<0>(b)); + time_point::duration(cuda::std::get<0>(a)) + time_point::duration(cuda::std::get<0>(b)); auto time_s = static_cast(time_d.count()) * static_cast(Period::num) / static_cast(Period::den); - T dist_km = thrust::get<1>(a) + thrust::get<1>(b); + T dist_km = cuda::std::get<1>(a) + cuda::std::get<1>(b); T dist_m = dist_km * T{1000.0}; // km to m T speed_m_s = dist_m / time_s; // m/ms to m/s return {time_d.count(), dist_km, dist_m, speed_m_s}; diff --git a/python/cuproj/cuproj/cuprojshim/include/detail/cuprojshim.hpp b/python/cuproj/cuproj/cuprojshim/include/detail/cuprojshim.hpp index 362d8cc43..2cca3d1d1 100644 --- a/python/cuproj/cuproj/cuprojshim/include/detail/cuprojshim.hpp +++ b/python/cuproj/cuproj/cuprojshim/include/detail/cuprojshim.hpp @@ -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. @@ -20,10 +20,10 @@ #include #include +#include #include #include #include -#include #include #include @@ -37,9 +37,8 @@ namespace detail { */ template > struct tuple_to_vec_2d { - __device__ VectorType operator()(thrust::tuple const& pos) - { - return VectorType{thrust::get<0>(pos), thrust::get<1>(pos)}; + __device__ VectorType operator()(cuda::std::tuple const &pos) { + return VectorType{cuda::std::get<0>(pos), cuda::std::get<1>(pos)}; } }; @@ -49,9 +48,8 @@ struct tuple_to_vec_2d { */ template > struct vec_2d_to_tuple { - __device__ thrust::tuple operator()(VectorType const& xy) - { - return thrust::make_tuple(xy.x, xy.y); + __device__ cuda::std::tuple operator()(VectorType const &xy) { + return cuda::std::make_tuple(xy.x, xy.y); } }; @@ -78,11 +76,11 @@ struct vec_2d_to_tuple { * "LegacyRandomAccessIterator" */ template -auto make_vec_2d_iterator(FirstIter first, SecondIter second) -{ +auto make_vec_2d_iterator(FirstIter first, SecondIter second) { using T = typename std::iterator_traits::value_type; - static_assert(std::is_same::value_type>(), - "Iterator value_type mismatch"); + static_assert( + std::is_same::value_type>(), + "Iterator value_type mismatch"); auto zipped = thrust::make_zip_iterator(first, second); return thrust::make_transform_iterator(zipped, tuple_to_vec_2d()); @@ -112,43 +110,33 @@ auto make_vec_2d_iterator(FirstIter first, SecondIter second) * "LegacyRandomAccessIterator" */ template -auto make_vec_2d_output_iterator(FirstIter first, SecondIter second) -{ - using T = typename std::iterator_traits::value_type; - auto zipped_out = thrust::make_zip_iterator(thrust::make_tuple(first, second)); +auto make_vec_2d_output_iterator(FirstIter first, SecondIter second) { + using T = typename std::iterator_traits::value_type; + auto zipped_out = thrust::make_zip_iterator(first, second); return thrust::transform_output_iterator(zipped_out, vec_2d_to_tuple()); } -} // namespace detail +} // namespace detail template -cuproj::projection>* make_projection(std::string const& src_epsg, - std::string const& dst_epsg) -{ +cuproj::projection> * +make_projection(std::string const &src_epsg, std::string const &dst_epsg) { return cuproj::make_projection>(src_epsg, dst_epsg); } template -void transform(cuproj::projection> const& proj, - cuproj::vec_2d* xy_in, - cuproj::vec_2d* xy_out, - std::size_t n, - cuproj::direction dir) -{ +void transform(cuproj::projection> const &proj, + cuproj::vec_2d *xy_in, cuproj::vec_2d *xy_out, + std::size_t n, cuproj::direction dir) { proj.transform(xy_in, xy_in + n, xy_out, dir); } template -void transform(cuproj::projection> const& proj, - T* x_in, - T* y_in, - T* x_out, - T* y_out, - std::size_t n, - cuproj::direction dir) -{ - auto xy_in = detail::make_vec_2d_iterator(x_in, y_in); +void transform(cuproj::projection> const &proj, T *x_in, + T *y_in, T *x_out, T *y_out, std::size_t n, + cuproj::direction dir) { + auto xy_in = detail::make_vec_2d_iterator(x_in, y_in); auto xy_out = detail::make_vec_2d_output_iterator(x_out, y_out); proj.transform(xy_in, xy_in + n, xy_out, dir); } -} // namespace cuprojshim +} // namespace cuprojshim