diff --git a/src/main/cpp/CMakeLists.txt b/src/main/cpp/CMakeLists.txt index fee3e60..7e26905 100644 --- a/src/main/cpp/CMakeLists.txt +++ b/src/main/cpp/CMakeLists.txt @@ -94,11 +94,8 @@ include(cmake/Modules/ConfigureCUDA.cmake) # set other CUDA compilation flags # ################################################################################################## # * dependencies ---------------------------------------------------------------------------------- -# find libcu++ -include(${rapids-cmake-dir}/cpm/libcudacxx.cmake) - -# find thrust/cub -include(${CUDF_DIR}/cpp/cmake/thirdparty/get_thrust.cmake) +# find CCCL +include(${CUDF_DIR}/cpp/cmake/thirdparty/get_cccl.cmake) # JNI find_package(JNI REQUIRED) diff --git a/src/main/cpp/src/map_utils.cu b/src/main/cpp/src/map_utils.cu index f8ac369..89c0d7c 100644 --- a/src/main/cpp/src/map_utils.cu +++ b/src/main/cpp/src/map_utils.cu @@ -51,6 +51,8 @@ #include #include +#include + // #include @@ -199,9 +201,12 @@ rmm::device_uvector compute_node_levels(int64_t num_nodes, }; auto const push_pop_it = thrust::make_transform_iterator( - tokens.begin(), [does_push, does_pop] __device__(PdaTokenT const token) -> cudf::size_type { + tokens.begin(), + cuda::proclaim_return_type( + [does_push, does_pop] __device__(PdaTokenT const token) -> cudf::size_type { return does_push(token) - does_pop(token); - }); + } + )); thrust::exclusive_scan( rmm::exec_policy(stream), push_pop_it, push_pop_it + tokens.size(), token_levels.begin()); diff --git a/src/main/cpp/src/row_conversion.cu b/src/main/cpp/src/row_conversion.cu index f2416fb..acec989 100644 --- a/src/main/cpp/src/row_conversion.cu +++ b/src/main/cpp/src/row_conversion.cu @@ -15,6 +15,7 @@ */ #include +#include #include #include #include @@ -1580,9 +1581,11 @@ batch_data build_batches(size_type num_rows, while (last_row_end < num_rows) { auto offset_row_sizes = thrust::make_transform_iterator( cumulative_row_sizes.begin(), - [last_row_end, cumulative_row_sizes = cumulative_row_sizes.data()] __device__(auto i) { - return i - cumulative_row_sizes[last_row_end]; - }); + cuda::proclaim_return_type( + [last_row_end, cumulative_row_sizes = cumulative_row_sizes.data()] __device__(auto i) { + return i - cumulative_row_sizes[last_row_end]; + }) + ); auto search_start = offset_row_sizes + last_row_end; auto search_end = offset_row_sizes + num_rows; @@ -1705,9 +1708,12 @@ size_type build_tiles( device_uvector tile_starts(num_batches + 1, stream); auto tile_iter = cudf::detail::make_counting_transform_iterator( - 0, [num_tiles = num_tiles.data(), num_batches] __device__(auto i) { - return (i < num_batches) ? num_tiles[i] : 0; - }); + 0, + cuda::proclaim_return_type( + [num_tiles = num_tiles.data(), num_batches] __device__(auto i) { + return (i < num_batches) ? num_tiles[i] : 0; + }) + ); thrust::exclusive_scan(rmm::exec_policy(stream), tile_iter, tile_iter + num_batches + 1, @@ -2459,7 +2465,9 @@ std::unique_ptr convert_from_rows(lists_column_view const& input, auto tmp = [num_rows, col_string_lengths] __device__(auto const& i) { return i < num_rows ? col_string_lengths[i] : 0; }; - auto bounded_iter = cudf::detail::make_counting_transform_iterator(0, tmp); + auto bounded_iter = + cudf::detail::make_counting_transform_iterator(0, + cuda::proclaim_return_type(tmp)); thrust::exclusive_scan(rmm::exec_policy(stream), bounded_iter, bounded_iter + num_rows + 1,