diff --git a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md index 74db9350f..325c9ce79 100644 --- a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md +++ b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md @@ -84,7 +84,7 @@ Examples: ```c++ template -void algorithm_function(int x, rmm::cuda_stream_view s, rmm::device_memory_resource* mr) +void algorithm_function(int x, rmm::cuda_stream_view s, rmm::device_async_resource_ref mr) { ... } @@ -233,9 +233,10 @@ std::unique_ptr points_in_spatial_window( cudf::column_view const& y); ``` -## RMM Memory Resources (`rmm::device_memory_resource`) +## Memory Resources (`rmm::device_memory_resource`) -libcuspatial allocates all device memory via RMM memory resources (MR). See the +libcuspatial allocates all device memory via RMM memory resources (MR) or CUDA MRs. Either type +can be passed to libcuspatial functions via `rmm::device_async_resource_ref` parameters. See the [RMM documentation](https://github.com/rapidsai/rmm/blob/main/README.md) for details. ### Current Device Memory Resource @@ -245,6 +246,27 @@ RMM provides a "default" memory resource for each device that can be accessed an respectively. All memory resource parameters should be defaulted to use the return value of `rmm::mr::get_current_device_resource()`. +### Resource Refs + +Memory resources are passed via resource ref parameters. A resource ref is memory resource wrapper +that enables consumers to specify properties of resources that they expect. These are defined +in the `cuda::mr` namespace of libcu++, but RMM provides some convenience wrappers in +`rmm/resource_ref.hpp`: + - `rmm::device_resource_ref` accepts a memory resource that provides synchronous allocation + of device-accessible memory. + - `rmm::device_async_resource_ref` accepts a memory resource that provides stream-ordered allocation + of device-accessible memory. + - `rmm::host_resource_ref` accepts a memory resource that provides synchronous allocation of host- + accessible memory. + - `rmm::host_async_resource_ref` accepts a memory resource that provides stream-ordered allocation + of host-accessible memory. + - `rmm::host_device_resource_ref` accepts a memory resource that provides synchronous allocation of + host- and device-accessible memory. + - `rmm::host_device_async_resource_ref` accepts a memory resource that provides stream-ordered + allocation of host- and device-accessible memory. + +See the libcu++ [docs on `resource_ref`](https://nvidia.github.io/cccl/libcudacxx/extended_api/memory_resource/resource_ref.html) for more information. + # libcuspatial API and Implementation This section provides specifics about the structure and implementation of cuSpatial API functions. @@ -439,8 +461,8 @@ There are a few key points to notice. cuSpatial APIs will not need to use this returned iterator. 9. All APIs that run CUDA device code (including Thrust algorithms) or allocate memory take a CUDA stream on which to execute the device code and allocate memory. - 10. Any API that allocate and return device data (not shown here) should also take an - `rmm::device_memory_resource` to use for output memory allocation. + 10. Any API that allocates and returns device data (not shown here) should also take an + `rmm::device_async_resource_ref` to use for output memory allocation. ### (Multiple) Return Values @@ -542,22 +564,28 @@ control how device memory is allocated. ### Output Memory -Any libcuspatial API that allocates memory that is *returned* to a user must accept a pointer to a -`device_memory_resource` as the last parameter. Inside the API, this memory resource must be used -to allocate any memory for returned objects. It should therefore be passed into functions whose -outputs will be returned. Example: +Any libcuspatial API that allocates memory that is *returned* to a user must accept a +`rmm::device_async_resource_ref` as the last parameter. Inside the API, this memory resource must +be used to allocate any memory for returned objects. It should therefore be passed into functions +whose outputs will be returned. Example: ```c++ // Returned `column` contains newly allocated memory, // therefore the API must accept a memory resource pointer std::unique_ptr returns_output_memory( - ..., rmm::device_memory_resource * mr = rmm::mr::get_current_device_resource()); + ..., rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); // This API does not allocate any new *output* memory, therefore // a memory resource is unnecessary void does_not_allocate_output_memory(...); ``` +This rule automatically applies to all detail APIs that allocate memory. Any detail API may be +called by any public API, and therefore could be allocating memory that is returned to the user. +To support such uses cases, all detail APIs allocating memory resources should accept an `mr` +parameter. Callers are responsible for either passing through a provided `mr` or +`rmm::mr::get_current_device_resource()` as needed. + ### Temporary Memory Not all memory allocated within a libcuspatial API is returned to the caller. Often algorithms must @@ -566,7 +594,7 @@ obtained from `rmm::mr::get_current_device_resource()` for temporary memory allo ```c++ rmm::device_buffer some_function( - ..., rmm::mr::device_memory_resource mr * = rmm::mr::get_current_device_resource()) { + ..., rmm::mr::device_async_resource_ref mr = rmm::mr::get_current_device_resource()) { rmm::device_buffer returned_buffer(..., mr); // Returned buffer uses the passed in MR ... rmm::device_buffer temporary_buffer(...); // Temporary buffer uses default MR @@ -578,12 +606,12 @@ rmm::device_buffer some_function( ### Memory Management libcuspatial code eschews raw pointers and direct memory allocation. Use RMM classes built to -use [`device_memory_resource`](https://github.com/rapidsai/rmm/#device_memory_resource) for device +use [memory resources](https://github.com/rapidsai/rmm/#device_memory_resource) for device memory allocation with automated lifetime management. #### rmm::device_buffer Allocates a specified number of bytes of untyped, uninitialized device memory using a -`device_memory_resource`. If no resource is explicitly provided, uses +memory resource. If no `rmm::device_async_resource_ref` is explicitly provided, uses `rmm::mr::get_current_device_resource()`. `rmm::device_buffer` is movable and copyable on a stream. A copy performs a deep copy of the diff --git a/cpp/include/cuspatial/bounding_boxes.hpp b/cpp/include/cuspatial/bounding_boxes.hpp index 9229eb7c2..2134bc1d1 100644 --- a/cpp/include/cuspatial/bounding_boxes.hpp +++ b/cpp/include/cuspatial/bounding_boxes.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,6 +19,7 @@ #include #include +#include #include @@ -53,7 +54,7 @@ std::unique_ptr linestring_bounding_boxes( cudf::column_view const& x, cudf::column_view const& y, double expansion_radius, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute minimum bounding box for each polygon in a list. @@ -80,8 +81,8 @@ std::unique_ptr polygon_bounding_boxes( cudf::column_view const& ring_offsets, cudf::column_view const& x, cudf::column_view const& y, - double expansion_radius = 0.0, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + double expansion_radius = 0.0, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/detail/index/construction/phase_1.cuh b/cpp/include/cuspatial/detail/index/construction/phase_1.cuh index 3d06c93b1..a444ae64a 100644 --- a/cpp/include/cuspatial/detail/index/construction/phase_1.cuh +++ b/cpp/include/cuspatial/detail/index/construction/phase_1.cuh @@ -24,6 +24,7 @@ #include #include #include +#include #include #include @@ -65,7 +66,7 @@ compute_point_keys_and_sorted_indices(PointIt points_first, T scale, int8_t max_depth, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_points = thrust::distance(points_first, points_last); rmm::device_uvector keys(num_points, stream); @@ -259,7 +260,7 @@ inline auto make_full_levels(PointIt points_first, T scale, int8_t max_depth, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_points = thrust::distance(points_first, points_last); // Compute point keys and sort into bottom-level quadrants diff --git a/cpp/include/cuspatial/detail/index/construction/phase_2.cuh b/cpp/include/cuspatial/detail/index/construction/phase_2.cuh index ffbb79c0d..88c4947e0 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-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -21,6 +21,7 @@ #include #include #include +#include #include #include @@ -309,7 +310,7 @@ inline rmm::device_uvector construct_non_leaf_indicator( int32_t num_parent_nodes, int32_t num_valid_nodes, int32_t max_size, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { // diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh index 1d88e5e36..818c6560b 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh @@ -34,6 +34,7 @@ #include #include #include +#include #include #include @@ -142,7 +143,7 @@ std::unique_ptr> compute_types_buffer( OffsetRangeB points_offset, OffsetRangeB segments_offset, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto types_buffer = std::make_unique>(union_column_size, stream, mr); thrust::tabulate(rmm::exec_policy(stream), @@ -162,7 +163,7 @@ std::unique_ptr> compute_types_buffer( template std::unique_ptr> compute_offset_buffer( rmm::device_uvector const& types_buffer, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { auto N = types_buffer.size(); @@ -202,7 +203,7 @@ template pairwise_linestring_intersection( MultiLinestringRange1 multilinestrings1, MultiLinestringRange2 multilinestrings2, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { using types_t = typename linestring_intersection_result::types_t; 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 e513cf651..f1a365555 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh @@ -29,6 +29,7 @@ #include #include #include +#include #include #include @@ -228,7 +229,7 @@ struct linestring_intersection_intermediates { /** @brief Construct a zero-pair, zero-geometry intermediate object */ linestring_intersection_intermediates(rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) : offsets(std::make_unique>(1, stream)), geoms(std::make_unique>(0, stream, mr)), lhs_linestring_ids(std::make_unique>(0, stream)), @@ -244,7 +245,7 @@ struct linestring_intersection_intermediates { std::size_t num_geoms, rmm::device_uvector const& num_geoms_per_pair, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) : offsets(std::make_unique>(num_pairs + 1, stream)), geoms(std::make_unique>(num_geoms, stream, mr)), lhs_linestring_ids(std::make_unique>(num_geoms, stream)), @@ -472,7 +473,7 @@ std::pair, index_t>, linestring_intersection_intermediates, index_t>> pairwise_linestring_intersection_with_duplicates(MultiLinestringRange1 multilinestrings1, MultiLinestringRange2 multilinestrings2, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { static_assert(std::is_integral_v, "Index type must be integral."); diff --git a/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh b/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh index 43ff7b0f6..472bb29b0 100644 --- a/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh +++ b/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh @@ -22,6 +22,8 @@ #include #include +#include + #include #include @@ -39,7 +41,7 @@ join_quadtree_and_bounding_boxes(point_quadtree_ref quadtree, T scale, int8_t max_depth, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { static_assert(is_same>(), "Iterator value_type mismatch"); 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 1ae0ff1bf..0b8f23e19 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, NVIDIA CORPORATION. + * Copyright (c) 2023-2024, 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,6 +26,7 @@ #include #include +#include #include #include @@ -108,7 +109,7 @@ std::pair, rmm::device_uvector> quadtr PointIterator points_first, MultiPolygonRange polygons, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { using T = iterator_vec_base_type; 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 6335f133f..5b61882b6 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 @@ -27,6 +27,7 @@ #include #include +#include #include #include @@ -156,7 +157,7 @@ quadtree_point_to_nearest_linestring(LinestringIndexIterator linestring_indices_ PointIterator points_first, MultiLinestringRange linestrings, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(linestrings.num_multilinestrings() == linestrings.num_linestrings(), "Only one linestring per multilinestring currently supported."); diff --git a/cpp/include/cuspatial/detail/point_quadtree.cuh b/cpp/include/cuspatial/detail/point_quadtree.cuh index 7e0d70dac..19922af35 100644 --- a/cpp/include/cuspatial/detail/point_quadtree.cuh +++ b/cpp/include/cuspatial/detail/point_quadtree.cuh @@ -24,6 +24,7 @@ #include #include #include +#include #include #include @@ -48,7 +49,7 @@ inline point_quadtree make_quad_tree(rmm::device_uvector& keys, int32_t max_size, int32_t level_1_size, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { // count the number of child nodes auto num_child_nodes = thrust::reduce(rmm::exec_policy(stream), @@ -155,7 +156,7 @@ inline point_quadtree make_leaf_tree(rmm::device_uvector& keys, rmm::device_uvector& lengths, int32_t num_top_quads, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { rmm::device_uvector levels(num_top_quads, stream, mr); rmm::device_uvector is_internal_node(num_top_quads, stream, mr); @@ -195,7 +196,7 @@ inline std::pair, point_quadtree> construct_quadtr T scale, int8_t max_depth, int32_t max_size, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { // Construct the full set of non-empty subquadrants starting from the lowest level. @@ -243,7 +244,7 @@ std::pair, point_quadtree> quadtree_on_points( int8_t max_depth, int32_t max_size, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_points = thrust::distance(points_first, points_last); if (num_points <= 0) { diff --git a/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh b/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh index 6da914181..11b8a594d 100644 --- a/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh +++ b/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh @@ -22,6 +22,7 @@ #include #include #include +#include #include #include @@ -63,7 +64,7 @@ void order_trajectories(IdInputIt ids_first, PointOutputIt points_out_first, TimestampOutputIt timestamps_out_first, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { using id_type = iterator_value_type; using timestamp_type = iterator_value_type; @@ -116,7 +117,7 @@ std::unique_ptr> derive_trajectories( PointOutputIt points_out_first, TimestampOutputIt timestamps_out_first, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { detail::order_trajectories(ids_first, ids_last, diff --git a/cpp/include/cuspatial/distance.hpp b/cpp/include/cuspatial/distance.hpp index 11cec4fae..76c920b11 100644 --- a/cpp/include/cuspatial/distance.hpp +++ b/cpp/include/cuspatial/distance.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,6 +25,7 @@ #include #include +#include #include @@ -53,8 +54,8 @@ std::unique_ptr haversine_distance( cudf::column_view const& a_lat, cudf::column_view const& b_lon, cudf::column_view const& b_lat, - double const radius = EARTH_RADIUS_KM, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + double const radius = EARTH_RADIUS_KM, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief computes Hausdorff distances for all pairs in a collection of spaces @@ -124,7 +125,7 @@ std::pair, cudf::table_view> directed_hausdorff_di cudf::column_view const& xs, cudf::column_view const& ys, cudf::column_view const& space_offsets, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)point-to-(multi)point Cartesian distance @@ -144,7 +145,7 @@ std::pair, cudf::table_view> directed_hausdorff_di std::unique_ptr pairwise_point_distance( geometry_column_view const& multipoints1, geometry_column_view const& multipoints2, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)points-to-(multi)linestrings Cartesian distance @@ -166,7 +167,7 @@ std::unique_ptr pairwise_point_distance( std::unique_ptr pairwise_point_linestring_distance( geometry_column_view const& multipoints, geometry_column_view const& multilinestrings, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)point-to-(multi)polygon Cartesian distance @@ -190,7 +191,7 @@ std::unique_ptr pairwise_point_linestring_distance( std::unique_ptr pairwise_point_polygon_distance( geometry_column_view const& multipoints, geometry_column_view const& multipolygons, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)linestring-to-(multi)linestring Cartesian distance @@ -212,7 +213,7 @@ std::unique_ptr pairwise_point_polygon_distance( std::unique_ptr pairwise_linestring_distance( geometry_column_view const& multilinestrings1, geometry_column_view const& multilinestrings2, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)linestring-to-(multi)polygon Cartesian distance @@ -237,7 +238,7 @@ std::unique_ptr pairwise_linestring_distance( std::unique_ptr pairwise_linestring_polygon_distance( geometry_column_view const& multilinestrings, geometry_column_view const& multipolygons, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)polygon-to-(multi)polygon Cartesian distance @@ -256,7 +257,7 @@ std::unique_ptr pairwise_linestring_polygon_distance( std::unique_ptr pairwise_polygon_distance( geometry_column_view const& multipolygons1, geometry_column_view const& multipolygons2, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/intersection.cuh b/cpp/include/cuspatial/intersection.cuh index 331e6eb0a..c086d84cc 100644 --- a/cpp/include/cuspatial/intersection.cuh +++ b/cpp/include/cuspatial/intersection.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, 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 @@ -87,8 +88,8 @@ template pairwise_linestring_intersection( MultiLinestringRange1 multilinestrings1, MultiLinestringRange2 multilinestrings2, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource(), - rmm::cuda_stream_view stream = rmm::cuda_stream_default); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource(), + rmm::cuda_stream_view stream = rmm::cuda_stream_default); } // namespace cuspatial diff --git a/cpp/include/cuspatial/intersection.hpp b/cpp/include/cuspatial/intersection.hpp index aedaeb521..1869768b1 100644 --- a/cpp/include/cuspatial/intersection.hpp +++ b/cpp/include/cuspatial/intersection.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * Copyright (c) 2023-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -21,6 +21,7 @@ #include #include +#include namespace cuspatial { /** @@ -54,6 +55,6 @@ struct linestring_intersection_column_result { linestring_intersection_column_result pairwise_linestring_intersection( geometry_column_view const& multilinestrings_lhs, geometry_column_view const& multilinestrings_rhs, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); } // namespace cuspatial diff --git a/cpp/include/cuspatial/nearest_points.hpp b/cpp/include/cuspatial/nearest_points.hpp index 640a1c9b6..46f0c14a2 100644 --- a/cpp/include/cuspatial/nearest_points.hpp +++ b/cpp/include/cuspatial/nearest_points.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,6 +19,8 @@ #include #include +#include + #include namespace cuspatial { @@ -165,7 +167,7 @@ point_linestring_nearest_points_result pairwise_point_linestring_nearest_points( std::optional> multilinestring_geometry_offsets, cudf::device_span linestring_part_offsets, cudf::column_view linestring_points_xy, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp b/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp index 9bfb13700..0df60c269 100644 --- a/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp +++ b/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp @@ -21,6 +21,7 @@ #include #include +#include #include @@ -69,6 +70,6 @@ namespace cuspatial { std::unique_ptr pairwise_multipoint_equals_count( geometry_column_view const& lhs, geometry_column_view const& rhs, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); } // namespace cuspatial diff --git a/cpp/include/cuspatial/point_in_polygon.hpp b/cpp/include/cuspatial/point_in_polygon.hpp index 11e7381c3..c66e2e260 100644 --- a/cpp/include/cuspatial/point_in_polygon.hpp +++ b/cpp/include/cuspatial/point_in_polygon.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -21,6 +21,7 @@ #include #include +#include #include @@ -78,7 +79,7 @@ std::unique_ptr point_in_polygon( cudf::column_view const& poly_ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Given (point, polygon pairs), tests whether the point of each pair is inside the polygon @@ -127,7 +128,7 @@ std::unique_ptr pairwise_point_in_polygon( cudf::column_view const& poly_ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/point_quadtree.cuh b/cpp/include/cuspatial/point_quadtree.cuh index a84b07558..f64cee6f0 100644 --- a/cpp/include/cuspatial/point_quadtree.cuh +++ b/cpp/include/cuspatial/point_quadtree.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, 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 @@ -175,8 +176,8 @@ std::pair, point_quadtree> quadtree_on_points( T scale, int8_t max_depth, int32_t max_size, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/point_quadtree.hpp b/cpp/include/cuspatial/point_quadtree.hpp index 17e156ea6..72306730f 100644 --- a/cpp/include/cuspatial/point_quadtree.hpp +++ b/cpp/include/cuspatial/point_quadtree.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,6 +19,7 @@ #include #include +#include #include @@ -74,7 +75,7 @@ std::pair, std::unique_ptr> quadtree_ double scale, int8_t max_depth, cudf::size_type max_size, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/points_in_range.hpp b/cpp/include/cuspatial/points_in_range.hpp index 7bd23c2c0..8637e2b24 100644 --- a/cpp/include/cuspatial/points_in_range.hpp +++ b/cpp/include/cuspatial/points_in_range.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -21,6 +21,7 @@ #include #include +#include #include @@ -60,7 +61,7 @@ std::unique_ptr points_in_range( double range_max_y, cudf::column_view const& x, cudf::column_view const& y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/projection.hpp b/cpp/include/cuspatial/projection.hpp index 1158d8c7c..e71164148 100644 --- a/cpp/include/cuspatial/projection.hpp +++ b/cpp/include/cuspatial/projection.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,6 +19,7 @@ #include #include +#include #include @@ -50,7 +51,7 @@ std::pair, std::unique_ptr> sinusoid double origin_lat, cudf::column_view const& input_lon, cudf::column_view const& input_lat, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/spatial_join.cuh b/cpp/include/cuspatial/spatial_join.cuh index 9815d5f4f..096af5fb2 100644 --- a/cpp/include/cuspatial/spatial_join.cuh +++ b/cpp/include/cuspatial/spatial_join.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, 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 @@ -69,8 +70,8 @@ join_quadtree_and_bounding_boxes( vec_2d const& v_min, T scale, int8_t max_depth, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Test whether the specified points are inside any of the specified polygons. @@ -127,8 +128,8 @@ std::pair, rmm::device_uvector> quadtr PointIndexIterator point_indices_last, PointIterator points_first, MultiPolygonRange polygons, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Finds the nearest linestring to each point in a quadrant, and computes the distances @@ -184,8 +185,8 @@ quadtree_point_to_nearest_linestring( PointIndexIterator point_indices_last, PointIterator points_first, MultiLinestringRange linestrings, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); } // namespace cuspatial diff --git a/cpp/include/cuspatial/spatial_join.hpp b/cpp/include/cuspatial/spatial_join.hpp index 579a3c4f1..e3bec92d8 100644 --- a/cpp/include/cuspatial/spatial_join.hpp +++ b/cpp/include/cuspatial/spatial_join.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,6 +19,7 @@ #include #include +#include #include @@ -71,7 +72,7 @@ std::unique_ptr join_quadtree_and_bounding_boxes( double y_max, double scale, int8_t max_depth, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Test whether the specified points are inside any of the specified polygons. @@ -122,7 +123,7 @@ std::unique_ptr quadtree_point_in_polygon( cudf::column_view const& ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Finds the nearest linestring to each point in a quadrant, and computes the distances @@ -171,7 +172,7 @@ std::unique_ptr quadtree_point_to_nearest_linestring( cudf::column_view const& linestring_offsets, cudf::column_view const& linestring_points_x, cudf::column_view const& linestring_points_y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/trajectory.cuh b/cpp/include/cuspatial/trajectory.cuh index f6c5d76f4..dc5ec9708 100644 --- a/cpp/include/cuspatial/trajectory.cuh +++ b/cpp/include/cuspatial/trajectory.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,6 +19,7 @@ #include #include #include +#include #include #include @@ -85,8 +86,8 @@ std::unique_ptr> derive_trajectories( IdOutputIt ids_output_first, PointOutputIt points_output_first, TimestampOutputIt timestamps_output_first, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute the total distance (in meters) and average speed (in m/s) of objects in diff --git a/cpp/include/cuspatial/trajectory.hpp b/cpp/include/cuspatial/trajectory.hpp index 5eb5eaf0d..308063569 100644 --- a/cpp/include/cuspatial/trajectory.hpp +++ b/cpp/include/cuspatial/trajectory.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, 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 @@ -62,7 +63,7 @@ std::pair, std::unique_ptr> derive_tr cudf::column_view const& x, cudf::column_view const& y, cudf::column_view const& timestamp, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute the distance and speed of objects in a trajectory. Groups the @@ -95,7 +96,7 @@ std::unique_ptr trajectory_distances_and_speeds( cudf::column_view const& x, cudf::column_view const& y, cudf::column_view const& timestamp, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute the spatial bounding boxes of trajectories. Groups the x, y, @@ -127,7 +128,7 @@ std::unique_ptr trajectory_bounding_boxes( cudf::column_view const& object_id, cudf::column_view const& x, cudf::column_view const& y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial_test/base_fixture.hpp b/cpp/include/cuspatial_test/base_fixture.hpp index c9927b489..820e54e77 100644 --- a/cpp/include/cuspatial_test/base_fixture.hpp +++ b/cpp/include/cuspatial_test/base_fixture.hpp @@ -20,6 +20,7 @@ #include #include +#include namespace cuspatial { namespace test { @@ -28,7 +29,7 @@ namespace test { * @brief Mixin to supply rmm resources for fixtures */ class RMMResourceMixin { - rmm::mr::device_memory_resource* _mr{rmm::mr::get_current_device_resource()}; + rmm::device_async_resource_ref _mr{rmm::mr::get_current_device_resource()}; rmm::cuda_stream_view _stream{rmm::cuda_stream_default}; public: @@ -37,7 +38,7 @@ class RMMResourceMixin { * all tests inheriting from this fixture * @return pointer to memory resource */ - rmm::mr::device_memory_resource* mr() { return _mr; } + rmm::device_async_resource_ref mr() { return _mr; } /** * @brief Returns `cuda_stream_view` that should be used for computation in diff --git a/cpp/include/cuspatial_test/vector_factories.cuh b/cpp/include/cuspatial_test/vector_factories.cuh index 139d411c5..a772cd6cf 100644 --- a/cpp/include/cuspatial_test/vector_factories.cuh +++ b/cpp/include/cuspatial_test/vector_factories.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, 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 @@ -56,7 +57,7 @@ auto make_device_vector(std::initializer_list inl) template auto make_device_uvector(std::initializer_list inl, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { std::vector hvec(inl.begin(), inl.end()); auto res = rmm::device_uvector(inl.size(), stream, mr); diff --git a/cpp/src/bounding_boxes/linestring_bounding_boxes.cu b/cpp/src/bounding_boxes/linestring_bounding_boxes.cu index 86b6aeb75..7716d231c 100644 --- a/cpp/src/bounding_boxes/linestring_bounding_boxes.cu +++ b/cpp/src/bounding_boxes/linestring_bounding_boxes.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, 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,6 +26,7 @@ #include #include +#include #include @@ -43,7 +44,7 @@ std::unique_ptr compute_linestring_bounding_boxes( cudf::column_view const& y, T expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_linestrings = linestring_offsets.size() > 0 ? linestring_offsets.size() - 1 : 0; @@ -93,7 +94,7 @@ struct dispatch_compute_linestring_bounding_boxes { cudf::column_view const& y, double expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return compute_linestring_bounding_boxes( linestring_offsets, x, y, static_cast(expansion_radius), stream, mr); @@ -109,7 +110,7 @@ std::unique_ptr linestring_bounding_boxes(cudf::column_view const& cudf::column_view const& y, double expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(x.type(), dispatch_compute_linestring_bounding_boxes{}, @@ -127,7 +128,7 @@ std::unique_ptr linestring_bounding_boxes(cudf::column_view const& cudf::column_view const& x, cudf::column_view const& y, double expansion_radius, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_linestrings = linestring_offsets.size() > 0 ? linestring_offsets.size() - 1 : 0; diff --git a/cpp/src/bounding_boxes/polygon_bounding_boxes.cu b/cpp/src/bounding_boxes/polygon_bounding_boxes.cu index e38e12dfd..c4c092e55 100644 --- a/cpp/src/bounding_boxes/polygon_bounding_boxes.cu +++ b/cpp/src/bounding_boxes/polygon_bounding_boxes.cu @@ -27,6 +27,7 @@ #include #include +#include #include @@ -43,7 +44,7 @@ std::unique_ptr compute_polygon_bounding_boxes(cudf::column_view co cudf::column_view const& y, T expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_polygons = poly_offsets.size() > 0 ? poly_offsets.size() - 1 : 0; @@ -96,7 +97,7 @@ struct dispatch_compute_polygon_bounding_boxes { cudf::column_view const& y, T expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return compute_polygon_bounding_boxes( poly_offsets, ring_offsets, x, y, expansion_radius, stream, mr); @@ -113,7 +114,7 @@ std::unique_ptr polygon_bounding_boxes(cudf::column_view const& pol cudf::column_view const& y, double expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(x.type(), dispatch_compute_polygon_bounding_boxes{}, @@ -133,7 +134,7 @@ std::unique_ptr polygon_bounding_boxes(cudf::column_view const& pol cudf::column_view const& x, cudf::column_view const& y, double expansion_radius, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_polys = poly_offsets.size() > 0 ? poly_offsets.size() - 1 : 0; auto num_rings = ring_offsets.size() > 0 ? ring_offsets.size() - 1 : 0; diff --git a/cpp/src/distance/hausdorff.cu b/cpp/src/distance/hausdorff.cu index 9706b5bf6..b4f2a62d1 100644 --- a/cpp/src/distance/hausdorff.cu +++ b/cpp/src/distance/hausdorff.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2023, NVIDIA CORPORATION. + * Copyright (c) 2019-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -32,6 +32,7 @@ #include #include +#include #include #include @@ -81,7 +82,7 @@ struct hausdorff_functor { cudf::column_view const& ys, cudf::column_view const& space_offsets, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto const num_points = static_cast(xs.size()); auto const num_spaces = static_cast(space_offsets.size()); @@ -120,7 +121,7 @@ std::pair, cudf::table_view> directed_hausdorff_di cudf::column_view const& xs, cudf::column_view const& ys, cudf::column_view const& space_offsets, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(xs.type() == ys.type(), "Inputs `xs` and `ys` must have same type."); CUSPATIAL_EXPECTS(xs.size() == ys.size(), "Inputs `xs` and `ys` must have same length."); diff --git a/cpp/src/distance/haversine.cu b/cpp/src/distance/haversine.cu index 3c4e631da..18a55cd2d 100644 --- a/cpp/src/distance/haversine.cu +++ b/cpp/src/distance/haversine.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, 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 @@ -49,7 +50,7 @@ struct haversine_functor { cudf::column_view const& b_lat, T radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { if (a_lon.is_empty()) { return cudf::empty_like(a_lon); } @@ -81,7 +82,7 @@ std::unique_ptr haversine_distance(cudf::column_view const& a_lon, cudf::column_view const& b_lat, double radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(radius > 0, "radius must be positive."); @@ -108,7 +109,7 @@ std::unique_ptr haversine_distance(cudf::column_view const& a_lon, cudf::column_view const& b_lon, cudf::column_view const& b_lat, double radius, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cuspatial::detail::haversine_distance( a_lon, a_lat, b_lon, b_lat, radius, rmm::cuda_stream_default, mr); diff --git a/cpp/src/distance/linestring_distance.cu b/cpp/src/distance/linestring_distance.cu index fe7b99fda..840d82b43 100644 --- a/cpp/src/distance/linestring_distance.cu +++ b/cpp/src/distance/linestring_distance.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, 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 @@ -42,7 +43,7 @@ struct pairwise_linestring_distance_launch { std::unique_ptr operator()(geometry_column_view const& multilinestrings1, geometry_column_view const& multilinestrings2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = multilinestrings1.size(); @@ -71,7 +72,7 @@ struct pairwise_linestring_distance_functor { std::unique_ptr operator()(geometry_column_view const& multilinestrings1, geometry_column_view const& multilinestrings2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(multilinestrings1.geometry_type() == geometry_type_id::LINESTRING && multilinestrings2.geometry_type() == geometry_type_id::LINESTRING, @@ -96,7 +97,7 @@ struct pairwise_linestring_distance_functor { std::unique_ptr pairwise_linestring_distance( geometry_column_view const& multilinestrings1, geometry_column_view const& multilinestrings2, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return multi_geometry_double_dispatch( multilinestrings1.collection_type(), diff --git a/cpp/src/distance/linestring_polygon_distance.cu b/cpp/src/distance/linestring_polygon_distance.cu index b6d370b90..552ab792d 100644 --- a/cpp/src/distance/linestring_polygon_distance.cu +++ b/cpp/src/distance/linestring_polygon_distance.cu @@ -34,6 +34,7 @@ #include #include +#include #include @@ -54,7 +55,7 @@ struct pairwise_linestring_polygon_distance_impl { std::unique_ptr operator()(geometry_column_view const& multilinestrings, geometry_column_view const& multipolygons, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto multilinestrings_range = make_multilinestring_range(multilinestrings); @@ -87,7 +88,7 @@ struct pairwise_linestring_polygon_distance { std::unique_ptr operator()(geometry_column_view const& multilinestrings, geometry_column_view const& multipolygons, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(multilinestrings.geometry_type() == geometry_type_id::LINESTRING && multipolygons.geometry_type() == geometry_type_id::POLYGON, @@ -114,7 +115,7 @@ struct pairwise_linestring_polygon_distance { std::unique_ptr pairwise_linestring_polygon_distance( geometry_column_view const& multilinestrings, geometry_column_view const& multipolygons, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return multi_geometry_double_dispatch( multilinestrings.collection_type(), diff --git a/cpp/src/distance/point_distance.cu b/cpp/src/distance/point_distance.cu index d2349b7a2..960c500d1 100644 --- a/cpp/src/distance/point_distance.cu +++ b/cpp/src/distance/point_distance.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, 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 @@ -40,7 +41,7 @@ struct pairwise_point_distance_impl { geometry_column_view const& multipoints1, geometry_column_view const& multipoints2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = multipoints1.size(); @@ -68,7 +69,7 @@ struct pairwise_point_distance_functor { std::unique_ptr operator()(geometry_column_view const& multipoints1, geometry_column_view const& multipoints2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(multipoints1.geometry_type() == geometry_type_id::POINT && multipoints2.geometry_type() == geometry_type_id::POINT, @@ -94,7 +95,7 @@ struct pairwise_point_distance_functor { std::unique_ptr pairwise_point_distance(geometry_column_view const& multipoints1, geometry_column_view const& multipoints2, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return multi_geometry_double_dispatch( multipoints1.collection_type(), diff --git a/cpp/src/distance/point_linestring_distance.cu b/cpp/src/distance/point_linestring_distance.cu index 59cb1b460..42d45b152 100644 --- a/cpp/src/distance/point_linestring_distance.cu +++ b/cpp/src/distance/point_linestring_distance.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -30,6 +30,7 @@ #include #include +#include #include #include @@ -48,7 +49,7 @@ struct pairwise_point_linestring_distance_impl { std::unique_ptr operator()(geometry_column_view const& multipoints, geometry_column_view const& multilinestrings, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = multipoints.size(); @@ -79,7 +80,7 @@ struct pairwise_point_linestring_distance_functor { std::unique_ptr operator()(geometry_column_view const& multipoints, geometry_column_view const& multilinestrings, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(multipoints.geometry_type() == geometry_type_id::POINT && multilinestrings.geometry_type() == geometry_type_id::LINESTRING, @@ -106,7 +107,7 @@ struct pairwise_point_linestring_distance_functor { std::unique_ptr pairwise_point_linestring_distance( geometry_column_view const& multipoints, geometry_column_view const& multilinestrings, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return multi_geometry_double_dispatch( multipoints.collection_type(), diff --git a/cpp/src/distance/point_polygon_distance.cu b/cpp/src/distance/point_polygon_distance.cu index dd7364fbd..b75b328bc 100644 --- a/cpp/src/distance/point_polygon_distance.cu +++ b/cpp/src/distance/point_polygon_distance.cu @@ -34,6 +34,7 @@ #include #include +#include #include @@ -54,7 +55,7 @@ struct pairwise_point_polygon_distance_impl { std::unique_ptr operator()(geometry_column_view const& multipoints, geometry_column_view const& multipolygons, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto multipoints_range = make_multipoint_range(multipoints); auto multipolygons_range = @@ -83,7 +84,7 @@ struct pairwise_point_polygon_distance { std::unique_ptr operator()(geometry_column_view const& multipoints, geometry_column_view const& multipolygons, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher( multipoints.coordinate_type(), @@ -100,7 +101,7 @@ struct pairwise_point_polygon_distance { std::unique_ptr pairwise_point_polygon_distance( geometry_column_view const& multipoints, geometry_column_view const& multipolygons, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(multipoints.geometry_type() == geometry_type_id::POINT && multipolygons.geometry_type() == geometry_type_id::POLYGON, diff --git a/cpp/src/distance/polygon_distance.cu b/cpp/src/distance/polygon_distance.cu index 678aa3f56..c5a6b0ed1 100644 --- a/cpp/src/distance/polygon_distance.cu +++ b/cpp/src/distance/polygon_distance.cu @@ -32,6 +32,7 @@ #include #include +#include #include @@ -52,7 +53,7 @@ struct pairwise_polygon_distance_impl { std::unique_ptr operator()(geometry_column_view const& lhs, geometry_column_view const& rhs, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto lhs_range = make_multipolygon_range(lhs); auto rhs_range = make_multipolygon_range(rhs); @@ -80,7 +81,7 @@ struct pairwise_polygon_distance { std::unique_ptr operator()(geometry_column_view const& lhs, geometry_column_view const& rhs, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(lhs.geometry_type() == geometry_type_id::POLYGON && rhs.geometry_type() == geometry_type_id::POLYGON, @@ -106,7 +107,7 @@ struct pairwise_polygon_distance { std::unique_ptr pairwise_polygon_distance(geometry_column_view const& lhs, geometry_column_view const& rhs, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return multi_geometry_double_dispatch( lhs.collection_type(), rhs.collection_type(), lhs, rhs, rmm::cuda_stream_default, mr); diff --git a/cpp/src/equality/pairwise_multipoint_equals_count.cu b/cpp/src/equality/pairwise_multipoint_equals_count.cu index 71e2b6d2b..abfde9e9c 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, NVIDIA CORPORATION. + * Copyright (c) 2023-2024, 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,6 +26,7 @@ #include #include +#include #include #include @@ -47,7 +48,7 @@ struct pairwise_multipoint_equals_count_impl { std::unique_ptr operator()(geometry_column_view const& lhs, geometry_column_view const& rhs, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = lhs.size(); // lhs is a buffer of xy coords auto type = cudf::data_type(cudf::type_to_id()); @@ -78,7 +79,7 @@ struct pairwise_multipoint_equals_count { std::unique_ptr operator()(geometry_column_view lhs, geometry_column_view rhs, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher( lhs.coordinate_type(), @@ -94,7 +95,7 @@ struct pairwise_multipoint_equals_count { std::unique_ptr pairwise_multipoint_equals_count(geometry_column_view const& lhs, geometry_column_view const& rhs, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(lhs.geometry_type() == geometry_type_id::POINT && rhs.geometry_type() == geometry_type_id::POINT, diff --git a/cpp/src/indexing/point_quadtree.cu b/cpp/src/indexing/point_quadtree.cu index 4e2f5987f..a8e53f1e1 100644 --- a/cpp/src/indexing/point_quadtree.cu +++ b/cpp/src/indexing/point_quadtree.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, 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,6 +26,7 @@ #include #include #include +#include #include @@ -72,7 +73,7 @@ struct dispatch_construct_quadtree { int8_t max_depth, cudf::size_type max_size, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto points = cuspatial::make_vec_2d_iterator(x.begin(), y.begin()); auto [point_indices, tree] = @@ -131,7 +132,7 @@ std::pair, std::unique_ptr> quadtree_ int8_t max_depth, cudf::size_type max_size, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(x.type(), dispatch_construct_quadtree{}, @@ -160,7 +161,7 @@ std::pair, std::unique_ptr> quadtree_ double scale, int8_t max_depth, cudf::size_type max_size, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(x.size() == y.size(), "x and y columns must have the same length"); if (x.is_empty() || y.is_empty()) { diff --git a/cpp/src/intersection/linestring_intersection.cu b/cpp/src/intersection/linestring_intersection.cu index 30e60e798..46734f69d 100644 --- a/cpp/src/intersection/linestring_intersection.cu +++ b/cpp/src/intersection/linestring_intersection.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * Copyright (c) 2023-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -34,6 +34,7 @@ #include #include +#include #include #include @@ -43,7 +44,7 @@ namespace detail { std::unique_ptr even_sequence(cudf::size_type size, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto res = cudf::make_numeric_column(cudf::data_type{cudf::type_to_id()}, size, @@ -67,7 +68,7 @@ struct pairwise_linestring_intersection_launch { operator()(geometry_column_view const& multilinestrings1, geometry_column_view const& multilinestrings2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { using index_t = cudf::size_type; @@ -155,7 +156,7 @@ struct pairwise_linestring_intersection { linestring_intersection_column_result operator()(geometry_column_view const& linestrings1, geometry_column_view const& linestrings2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(linestrings1.coordinate_type() == linestrings2.coordinate_type(), "Input linestring coordinates must be the same type."); @@ -176,7 +177,7 @@ struct pairwise_linestring_intersection { linestring_intersection_column_result pairwise_linestring_intersection( geometry_column_view const& lhs, geometry_column_view const& rhs, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(lhs.geometry_type() == geometry_type_id::LINESTRING && rhs.geometry_type() == geometry_type_id::LINESTRING, diff --git a/cpp/src/join/quadtree_bbox_filtering.cu b/cpp/src/join/quadtree_bbox_filtering.cu index 59bd369cd..1757938f3 100644 --- a/cpp/src/join/quadtree_bbox_filtering.cu +++ b/cpp/src/join/quadtree_bbox_filtering.cu @@ -27,6 +27,7 @@ #include #include #include +#include #include @@ -42,7 +43,7 @@ struct dispatch_quadtree_bounding_box_join { double y_min, double scale, int8_t max_depth, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { auto bbox_min = cuspatial::make_vec_2d_iterator(bbox.column(0).template begin(), @@ -94,7 +95,7 @@ std::unique_ptr join_quadtree_and_bounding_boxes(cudf::table_view c double y_max, double scale, int8_t max_depth, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(quadtree.num_columns() == 5, "quadtree table must have 5 columns"); CUSPATIAL_EXPECTS(bbox.num_columns() == 4, "bbox table must have 4 columns"); diff --git a/cpp/src/join/quadtree_point_in_polygon.cu b/cpp/src/join/quadtree_point_in_polygon.cu index 75b723aa8..48f9055ba 100644 --- a/cpp/src/join/quadtree_point_in_polygon.cu +++ b/cpp/src/join/quadtree_point_in_polygon.cu @@ -28,6 +28,7 @@ #include #include +#include namespace cuspatial { namespace detail { @@ -52,7 +53,7 @@ struct compute_quadtree_point_in_polygon { cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto poly_indices = poly_quad_pairs.column(0); auto quad_indices = poly_quad_pairs.column(1); @@ -120,7 +121,7 @@ std::unique_ptr quadtree_point_in_polygon(cudf::table_view const& p cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(point_x.type(), compute_quadtree_point_in_polygon{}, @@ -148,7 +149,7 @@ std::unique_ptr quadtree_point_in_polygon(cudf::table_view const& p cudf::column_view const& ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(poly_quad_pairs.num_columns() == 2, "a quadrant-polygon table must have 2 columns"); diff --git a/cpp/src/join/quadtree_point_to_nearest_linestring.cu b/cpp/src/join/quadtree_point_to_nearest_linestring.cu index c423c1cb4..8288ecbc3 100644 --- a/cpp/src/join/quadtree_point_to_nearest_linestring.cu +++ b/cpp/src/join/quadtree_point_to_nearest_linestring.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, 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 @@ -54,7 +55,7 @@ struct compute_quadtree_point_to_nearest_linestring { cudf::column_view const& linestring_points_x, cudf::column_view const& linestring_points_y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto linestring_indices = linestring_quad_pairs.column(0); auto quad_indices = linestring_quad_pairs.column(1); @@ -120,7 +121,7 @@ std::unique_ptr quadtree_point_to_nearest_linestring( cudf::column_view const& linestring_points_x, cudf::column_view const& linestring_points_y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(point_x.type(), compute_quadtree_point_to_nearest_linestring{}, @@ -147,7 +148,7 @@ std::unique_ptr quadtree_point_to_nearest_linestring( cudf::column_view const& linestring_offsets, cudf::column_view const& linestring_points_x, cudf::column_view const& linestring_points_y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(linestring_quad_pairs.num_columns() == 2, "a quadrant-linestring table must have 2 columns"); diff --git a/cpp/src/nearest_points/point_linestring_nearest_points.cu b/cpp/src/nearest_points/point_linestring_nearest_points.cu index 9e15cc225..00a4d365f 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-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -30,6 +30,7 @@ #include #include +#include #include #include @@ -56,7 +57,7 @@ struct pairwise_point_linestring_nearest_points_impl { cudf::device_span linestring_offsets, cudf::column_view linestring_points_xy, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_points = static_cast(points_xy.size() / 2); auto num_linestring_points = static_cast(linestring_points_xy.size() / 2); @@ -219,7 +220,7 @@ struct pairwise_point_linestring_nearest_points_functor { cudf::device_span linestring_part_offsets, cudf::column_view linestring_points_xy, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(points_xy.size() % 2 == 0 && linestring_points_xy.size() % 2 == 0, "Points array must contain even number of coordinates."); @@ -258,7 +259,7 @@ point_linestring_nearest_points_result pairwise_point_linestring_nearest_points( std::optional> multilinestring_geometry_offsets, cudf::device_span linestring_part_offsets, cudf::column_view linestring_points_xy, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return double_boolean_dispatch( multipoint_geometry_offsets.has_value(), diff --git a/cpp/src/point_in_polygon/point_in_polygon.cu b/cpp/src/point_in_polygon/point_in_polygon.cu index e6eb56d08..c22a8da94 100644 --- a/cpp/src/point_in_polygon/point_in_polygon.cu +++ b/cpp/src/point_in_polygon/point_in_polygon.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, 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 @@ -56,7 +57,7 @@ struct point_in_polygon_functor { cudf::column_view const& poly_points_y, bool pairwise, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = test_points_x.size(); auto tid = pairwise ? cudf::type_to_id() : cudf::type_to_id(); @@ -112,7 +113,7 @@ std::unique_ptr point_in_polygon(cudf::column_view const& test_poi cudf::column_view const& poly_points_y, bool pairwise, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS( test_points_x.size() == test_points_y.size() and poly_points_x.size() == poly_points_y.size(), @@ -155,7 +156,7 @@ std::unique_ptr point_in_polygon(cudf::column_view const& test_poi cudf::column_view const& poly_ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cuspatial::detail::point_in_polygon(test_points_x, test_points_y, @@ -174,7 +175,7 @@ std::unique_ptr pairwise_point_in_polygon(cudf::column_view const& cudf::column_view const& poly_ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cuspatial::detail::point_in_polygon(test_points_x, test_points_y, diff --git a/cpp/src/points_in_range/points_in_range.cu b/cpp/src/points_in_range/points_in_range.cu index df6f60c6a..edb843a93 100644 --- a/cpp/src/points_in_range/points_in_range.cu +++ b/cpp/src/points_in_range/points_in_range.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,6 +25,7 @@ #include #include +#include #include #include @@ -42,7 +43,7 @@ struct points_in_range_dispatch { cudf::column_view const& x, cudf::column_view const& y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto points_begin = cuspatial::make_vec_2d_iterator(x.begin(), y.begin()); @@ -98,7 +99,7 @@ std::unique_ptr points_in_range(double range_min_x, cudf::column_view const& x, cudf::column_view const& y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(x.type() == y.type(), "Type mismatch between x and y arrays"); CUSPATIAL_EXPECTS(x.size() == y.size(), "Size mismatch between x and y arrays"); @@ -129,7 +130,7 @@ std::unique_ptr points_in_range(double range_min_x, double range_max_y, cudf::column_view const& x, cudf::column_view const& y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return detail::points_in_range( range_min_x, range_max_x, range_min_y, range_max_y, x, y, rmm::cuda_stream_default, mr); diff --git a/cpp/src/projection/sinusoidal_projection.cu b/cpp/src/projection/sinusoidal_projection.cu index 1be4ade40..2552601a4 100644 --- a/cpp/src/projection/sinusoidal_projection.cu +++ b/cpp/src/projection/sinusoidal_projection.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2023, NVIDIA CORPORATION. + * Copyright (c) 2019-2024, 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,6 +26,7 @@ #include #include +#include #include #include @@ -53,7 +54,7 @@ struct dispatch_sinusoidal_projection { cudf::column_view const& input_lon, cudf::column_view const& input_lat, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = input_lon.size(); auto type = cudf::data_type{cudf::type_to_id()}; @@ -87,7 +88,7 @@ pair_of_columns sinusoidal_projection(double origin_lon, cudf::column_view const& input_lon, cudf::column_view const& input_lat, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS( origin_lon >= -180 && origin_lon <= 180 && origin_lat >= -90 && origin_lat <= 90, @@ -116,7 +117,7 @@ pair_of_columns sinusoidal_projection(double origin_lon, double origin_lat, cudf::column_view const& input_lon, cudf::column_view const& input_lat, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return detail::sinusoidal_projection( origin_lon, origin_lat, input_lon, input_lat, rmm::cuda_stream_default, mr); diff --git a/cpp/src/trajectory/derive_trajectories.cu b/cpp/src/trajectory/derive_trajectories.cu index 8356369ef..e4e6f3745 100644 --- a/cpp/src/trajectory/derive_trajectories.cu +++ b/cpp/src/trajectory/derive_trajectories.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,6 +25,7 @@ #include #include +#include #include #include @@ -43,7 +44,7 @@ struct derive_trajectories_dispatch { cudf::column_view const& y, cudf::column_view const& timestamp, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto cols = std::vector>{}; cols.reserve(4); @@ -95,7 +96,7 @@ std::pair, std::unique_ptr> derive_tr cudf::column_view const& y, cudf::column_view const& timestamp, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::double_type_dispatcher(x.type(), timestamp.type(), @@ -114,7 +115,7 @@ std::pair, std::unique_ptr> derive_tr cudf::column_view const& x, cudf::column_view const& y, cudf::column_view const& timestamp, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS( x.size() == y.size() && x.size() == object_id.size() && x.size() == timestamp.size(), diff --git a/cpp/src/trajectory/trajectory_bounding_boxes.cu b/cpp/src/trajectory/trajectory_bounding_boxes.cu index 8d441b9c4..baeea23e4 100644 --- a/cpp/src/trajectory/trajectory_bounding_boxes.cu +++ b/cpp/src/trajectory/trajectory_bounding_boxes.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, 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 @@ -45,7 +46,7 @@ struct dispatch_element { cudf::column_view const& x, cudf::column_view const& y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { // Construct output columns auto type = cudf::data_type{cudf::type_to_id()}; @@ -101,7 +102,7 @@ std::unique_ptr trajectory_bounding_boxes(cudf::size_type num_traje cudf::column_view const& x, cudf::column_view const& y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher( x.type(), dispatch_element{}, num_trajectories, object_id, x, y, stream, mr); @@ -112,7 +113,7 @@ std::unique_ptr trajectory_bounding_boxes(cudf::size_type num_traje cudf::column_view const& object_id, cudf::column_view const& x, cudf::column_view const& y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(object_id.size() == x.size() && x.size() == y.size(), "Data size mismatch"); CUSPATIAL_EXPECTS(x.type().id() == y.type().id(), "Data type mismatch"); diff --git a/cpp/src/trajectory/trajectory_distances_and_speeds.cu b/cpp/src/trajectory/trajectory_distances_and_speeds.cu index a27f8b7f8..89474e925 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, NVIDIA CORPORATION. + * Copyright (c) 2020-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -29,6 +29,7 @@ #include #include +#include #include #include @@ -54,7 +55,7 @@ struct dispatch_timestamp { cudf::column_view const& y, cudf::column_view const& timestamp, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { // Construct output columns auto type = cudf::data_type{cudf::type_to_id()}; @@ -102,7 +103,7 @@ struct dispatch_element { cudf::column_view const& y, cudf::column_view const& timestamp, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(timestamp.type(), dispatch_timestamp{}, @@ -132,7 +133,7 @@ std::unique_ptr trajectory_distances_and_speeds(cudf::size_type num cudf::column_view const& y, cudf::column_view const& timestamp, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher( x.type(), dispatch_element{}, num_trajectories, object_id, x, y, timestamp, stream, mr); @@ -144,7 +145,7 @@ std::unique_ptr trajectory_distances_and_speeds(cudf::size_type num cudf::column_view const& x, cudf::column_view const& y, cudf::column_view const& timestamp, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS( x.size() == y.size() && x.size() == object_id.size() && x.size() == timestamp.size(), diff --git a/cpp/tests/distance/linestring_polygon_distance_test.cu b/cpp/tests/distance/linestring_polygon_distance_test.cu index 7cc3b6b55..9a09fd84b 100644 --- a/cpp/tests/distance/linestring_polygon_distance_test.cu +++ b/cpp/tests/distance/linestring_polygon_distance_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * Copyright (c) 2023-2024, 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,6 +26,7 @@ #include #include #include +#include #include @@ -35,7 +36,7 @@ using namespace cuspatial::test; template struct PairwiseLinestringPolygonDistanceTest : public BaseFixture { rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; } - rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); } + rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); } void run_single(std::initializer_list multilinestring_geometry_offsets, std::initializer_list multilinestring_part_offsets, diff --git a/cpp/tests/distance/point_polygon_distance_test.cu b/cpp/tests/distance/point_polygon_distance_test.cu index c16d324e0..244ac5ee6 100644 --- a/cpp/tests/distance/point_polygon_distance_test.cu +++ b/cpp/tests/distance/point_polygon_distance_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * Copyright (c) 2023-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,6 +25,7 @@ #include #include #include +#include #include @@ -36,7 +37,7 @@ double constexpr PI = 3.14159265358979323846; template struct PairwisePointPolygonDistanceTest : public ::testing::Test { rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; } - rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); } + rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); } void run_single(std::initializer_list>> multipoints, std::initializer_list multipolygon_geometry_offsets, diff --git a/cpp/tests/intersection/intersection_test_utils.cuh b/cpp/tests/intersection/intersection_test_utils.cuh index ab0971648..3a2112d07 100644 --- a/cpp/tests/intersection/intersection_test_utils.cuh +++ b/cpp/tests/intersection/intersection_test_utils.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, 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,8 @@ #include #include +#include + #include #include #include @@ -89,7 +91,7 @@ struct order_key_value_pairs { template linestring_intersection_result segment_sort_intersection_result( linestring_intersection_result& result, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { auto const num_points = result.points_coords->size(); @@ -189,7 +191,7 @@ auto make_linestring_intersection_result( std::initializer_list rhs_linestring_ids, std::initializer_list rhs_segment_ids, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto d_geometry_collection_offset = make_device_uvector(geometry_collection_offset, stream, mr); diff --git a/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu b/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu index eb36d5e35..760ab72a4 100644 --- a/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu +++ b/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-2024, 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 @@ -44,7 +45,7 @@ auto make_linestring_intersection_intermediates(std::initializer_list std::initializer_list rhs_linestring_ids, std::initializer_list rhs_segment_ids, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto d_offset = make_device_uvector(offsets, stream, mr); auto d_geoms = make_device_uvector(geoms, stream, mr); @@ -66,7 +67,7 @@ auto make_linestring_intersection_intermediates(std::initializer_list template struct LinestringIntersectionIntermediatesRemoveIfTest : public ::testing::Test { rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; } - rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); } + rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); } template void run_single(IntermediateType& intermediates, diff --git a/cpp/tests/intersection/linestring_intersection_test.cpp b/cpp/tests/intersection/linestring_intersection_test.cpp index 9b9318104..d91d99ddb 100644 --- a/cpp/tests/intersection/linestring_intersection_test.cpp +++ b/cpp/tests/intersection/linestring_intersection_test.cpp @@ -37,6 +37,7 @@ #include #include +#include #include #include @@ -116,7 +117,7 @@ std::pair> make_linestring_col struct LinestringIntersectionTestBase : public BaseFixture { rmm::cuda_stream_view stream{rmm::cuda_stream_default}; - rmm::mr::device_memory_resource* mr{rmm::mr::get_current_device_resource()}; + rmm::device_async_resource_ref mr{rmm::mr::get_current_device_resource()}; }; template diff --git a/cpp/tests/intersection/linestring_intersection_test.cu b/cpp/tests/intersection/linestring_intersection_test.cu index 7f758821c..d6ec4affd 100644 --- a/cpp/tests/intersection/linestring_intersection_test.cu +++ b/cpp/tests/intersection/linestring_intersection_test.cu @@ -30,6 +30,7 @@ #include #include +#include #include #include @@ -43,7 +44,7 @@ using namespace cuspatial::test; template struct LinestringIntersectionTest : public ::testing::Test { rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; } - rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); } + rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); } template void run_single_test(MultiLinestringRange lhs, diff --git a/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu b/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu index 359b45700..720eb2efa 100644 --- a/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu +++ b/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu @@ -31,6 +31,7 @@ #include #include #include +#include #include #include @@ -72,7 +73,7 @@ using namespace cuspatial::test; template Intermediate segmented_sort_intersection_intermediates(Intermediate& intermediate, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { using GeomType = typename Intermediate::geometry_t; using IndexType = typename Intermediate::index_t; @@ -126,7 +127,7 @@ Intermediate segmented_sort_intersection_intermediates(Intermediate& intermediat template struct LinestringIntersectionDuplicatesTest : public ::testing::Test { rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; } - rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); } + rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); } template void run_single(MultilinestringRange1 lhs, diff --git a/cpp/tests/join/quadtree_point_in_polygon_test_large.cu b/cpp/tests/join/quadtree_point_in_polygon_test_large.cu index d80709486..e1f0c86f6 100644 --- a/cpp/tests/join/quadtree_point_in_polygon_test_large.cu +++ b/cpp/tests/join/quadtree_point_in_polygon_test_large.cu @@ -31,6 +31,7 @@ #include #include #include +#include #include #include @@ -54,7 +55,7 @@ inline auto generate_points( uint32_t points_per_quad, std::size_t seed, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()) { auto engine = cuspatial::test::deterministic_engine(0); auto uniform = cuspatial::test::make_normal_dist(0.0, 1.0);