From fea4a00e9c229eb4470fa643d3f66cdb8b2c76c6 Mon Sep 17 00:00:00 2001 From: Jason Lowe Date: Mon, 20 Mar 2023 20:06:00 -0500 Subject: [PATCH] Upmerge cudf and fix API changes to make_device_uvector_async (#1012) Signed-off-by: Jason Lowe --- src/main/cpp/src/map_utils.cu | 3 +- src/main/cpp/src/row_conversion.cu | 64 ++++++++++++++++++++---------- thirdparty/cudf | 2 +- 3 files changed, 45 insertions(+), 24 deletions(-) diff --git a/src/main/cpp/src/map_utils.cu b/src/main/cpp/src/map_utils.cu index 590fd04cc4..1e6e33fe8c 100644 --- a/src/main/cpp/src/map_utils.cu +++ b/src/main/cpp/src/map_utils.cu @@ -67,7 +67,8 @@ namespace { rmm::device_uvector unify_json_strings(cudf::column_view const &input, rmm::cuda_stream_view stream) { if (input.is_empty()) { - return cudf::detail::make_device_uvector_async(std::vector{'[', ']'}, stream); + return cudf::detail::make_device_uvector_async(std::vector{'[', ']'}, stream, + rmm::mr::get_current_device_resource()); } auto const d_strings = cudf::column_device_view::create(input, stream); diff --git a/src/main/cpp/src/row_conversion.cu b/src/main/cpp/src/row_conversion.cu index 6c50414784..dc45220091 100644 --- a/src/main/cpp/src/row_conversion.cu +++ b/src/main/cpp/src/row_conversion.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -225,7 +225,8 @@ build_string_row_offsets(table_view const &tbl, size_type fixed_width_and_validi std::copy_if(offsets_iter, offsets_iter + tbl.num_columns(), std::back_inserter(offsets_iterators), [](auto const &offset_ptr) { return offset_ptr != nullptr; }); - return make_device_uvector_async(offsets_iterators, stream); + return make_device_uvector_async(offsets_iterators, stream, + rmm::mr::get_current_device_resource()); }(); auto const num_columns = static_cast(d_offsets_iterators.size()); @@ -1538,7 +1539,9 @@ batch_data build_batches(size_type num_rows, RowSize row_sizes, bool all_fixed_w last_row_end = row_end; } - return {std::move(batch_row_offsets), make_device_uvector_async(batch_row_boundaries, stream), + return {std::move(batch_row_offsets), + make_device_uvector_async(batch_row_boundaries, stream, + rmm::mr::get_current_device_resource()), std::move(batch_row_boundaries), std::move(row_batches)}; } @@ -1749,8 +1752,10 @@ std::vector> convert_to_rows( return table_view(cols); }; - auto dev_col_sizes = make_device_uvector_async(column_info.column_sizes, stream); - auto dev_col_starts = make_device_uvector_async(column_info.column_starts, stream); + auto dev_col_sizes = make_device_uvector_async(column_info.column_sizes, stream, + rmm::mr::get_current_device_resource()); + auto dev_col_starts = make_device_uvector_async(column_info.column_starts, stream, + rmm::mr::get_current_device_resource()); // Get the pointers to the input columnar data ready auto const data_begin = thrust::make_transform_iterator(tbl.begin(), [](auto const &c) { @@ -1763,8 +1768,10 @@ std::vector> convert_to_rows( thrust::make_transform_iterator(tbl.begin(), [](auto const &c) { return c.null_mask(); }); std::vector input_nm(nm_begin, nm_begin + tbl.num_columns()); - auto dev_input_data = make_device_uvector_async(input_data, stream); - auto dev_input_nm = make_device_uvector_async(input_nm, stream); + auto dev_input_data = + make_device_uvector_async(input_data, stream, rmm::mr::get_current_device_resource()); + auto dev_input_nm = + make_device_uvector_async(input_nm, stream, rmm::mr::get_current_device_resource()); // the first batch always exists unless we were sent an empty table auto const first_batch_size = batch_info.row_batches[0].row_count; @@ -1810,7 +1817,8 @@ std::vector> convert_to_rows( auto validity_tile_infos = detail::build_validity_tile_infos( tbl.num_columns(), num_rows, shmem_limit_per_tile, batch_info.row_batches); - auto dev_validity_tile_infos = make_device_uvector_async(validity_tile_infos, stream); + auto dev_validity_tile_infos = make_device_uvector_async(validity_tile_infos, stream, + rmm::mr::get_current_device_resource()); auto const validity_offset = column_info.column_starts.back(); @@ -1846,9 +1854,10 @@ std::vector> convert_to_rows( std::vector variable_width_input_data( variable_data_begin, variable_data_begin + variable_width_table.num_columns()); - auto dev_variable_input_data = make_device_uvector_async(variable_width_input_data, stream); - auto dev_variable_col_output_offsets = - make_device_uvector_async(column_info.variable_width_column_starts, stream); + auto dev_variable_input_data = make_device_uvector_async( + variable_width_input_data, stream, rmm::mr::get_current_device_resource()); + auto dev_variable_col_output_offsets = make_device_uvector_async( + column_info.variable_width_column_starts, stream, rmm::mr::get_current_device_resource()); for (uint i = 0; i < batch_info.row_batches.size(); i++) { auto const batch_row_offset = batch_info.batch_row_boundaries[i]; @@ -2075,8 +2084,10 @@ std::unique_ptr convert_from_rows(lists_column_view const &input, // Ideally we would check that the offsets are all the same, etc. but for now this is probably // fine CUDF_EXPECTS(size_per_row * num_rows <= child.size(), "The layout of the data appears to be off"); - auto dev_col_starts = make_device_uvector_async(column_info.column_starts, stream); - auto dev_col_sizes = make_device_uvector_async(column_info.column_sizes, stream); + auto dev_col_starts = make_device_uvector_async(column_info.column_starts, stream, + rmm::mr::get_current_device_resource()); + auto dev_col_sizes = make_device_uvector_async(column_info.column_sizes, stream, + rmm::mr::get_current_device_resource()); // Allocate the columns we are going to write into std::vector> output_columns; @@ -2117,16 +2128,20 @@ std::unique_ptr
convert_from_rows(lists_column_view const &input, } } - auto dev_string_row_offsets = make_device_uvector_async(string_row_offsets, stream); - auto dev_string_lengths = make_device_uvector_async(string_lengths, stream); + auto dev_string_row_offsets = + make_device_uvector_async(string_row_offsets, stream, rmm::mr::get_current_device_resource()); + auto dev_string_lengths = + make_device_uvector_async(string_lengths, stream, rmm::mr::get_current_device_resource()); // build the row_batches from the passed in list column std::vector row_batches; row_batches.push_back( {detail::row_batch{child.size(), num_rows, device_uvector(0, stream)}}); - auto dev_output_data = make_device_uvector_async(output_data, stream); - auto dev_output_nm = make_device_uvector_async(output_nm, stream); + auto dev_output_data = + make_device_uvector_async(output_data, stream, rmm::mr::get_current_device_resource()); + auto dev_output_nm = + make_device_uvector_async(output_nm, stream, rmm::mr::get_current_device_resource()); // only ever get a single batch when going from rows, so boundaries are 0, num_rows constexpr auto num_batches = 2; @@ -2163,7 +2178,8 @@ std::unique_ptr
convert_from_rows(lists_column_view const &input, auto validity_tile_infos = detail::build_validity_tile_infos(schema.size(), num_rows, shmem_limit_per_tile, row_batches); - auto dev_validity_tile_infos = make_device_uvector_async(validity_tile_infos, stream); + auto dev_validity_tile_infos = make_device_uvector_async(validity_tile_infos, stream, + rmm::mr::get_current_device_resource()); dim3 const validity_blocks(validity_tile_infos.size()); @@ -2220,8 +2236,10 @@ std::unique_ptr
convert_from_rows(lists_column_view const &input, string_col_offsets.push_back(std::move(output_string_offsets)); string_data_cols.push_back(std::move(string_data)); } - auto dev_string_col_offsets = make_device_uvector_async(string_col_offset_ptrs, stream); - auto dev_string_data_cols = make_device_uvector_async(string_data_col_ptrs, stream); + auto dev_string_col_offsets = make_device_uvector_async(string_col_offset_ptrs, stream, + rmm::mr::get_current_device_resource()); + auto dev_string_data_cols = make_device_uvector_async(string_data_col_ptrs, stream, + rmm::mr::get_current_device_resource()); dim3 const string_blocks( std::min(std::max(MIN_STRING_BLOCKS, num_rows / NUM_STRING_ROWS_PER_BLOCK_FROM_ROWS), @@ -2273,8 +2291,10 @@ std::unique_ptr
convert_from_rows_fixed_width_optimized( // fine CUDF_EXPECTS(size_per_row * num_rows == child.size(), "The layout of the data appears to be off"); - auto dev_column_start = make_device_uvector_async(column_start, stream); - auto dev_column_size = make_device_uvector_async(column_size, stream); + auto dev_column_start = + make_device_uvector_async(column_start, stream, rmm::mr::get_current_device_resource()); + auto dev_column_size = + make_device_uvector_async(column_size, stream, rmm::mr::get_current_device_resource()); // Allocate the columns we are going to write into std::vector> output_columns; diff --git a/thirdparty/cudf b/thirdparty/cudf index 8fbfb4ae5f..3b8064de10 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 8fbfb4ae5f083ac8473023f9e4982e8021b83c6d +Subproject commit 3b8064de10e1891c7d104a6e63bb68393a15fc68