Skip to content

Commit

Permalink
Upmerge cudf and fix API changes to make_device_uvector_async (#1012)
Browse files Browse the repository at this point in the history 
Signed-off-by: Jason Lowe <[email protected]>
  • Loading branch information
@jlowe
jlowe authored Mar 21, 2023
1 parent 4f122d7 commit fea4a00
Show file tree
Hide file tree
Showing 3 changed files with 45 additions and 24 deletions.
3 changes: 2 additions & 1 deletion src/main/cpp/src/map_utils.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -67,7 +67,8 @@ namespace {
rmm::device_uvector<char> unify_json_strings(cudf::column_view const &input,
rmm::cuda_stream_view stream) {
if (input.is_empty()) {
return cudf::detail::make_device_uvector_async<char>(std::vector<char>{'[', ']'}, stream);
return cudf::detail::make_device_uvector_async<char>(std::vector<char>{'[', ']'}, stream,
rmm::mr::get_current_device_resource());
}

auto const d_strings = cudf::column_device_view::create(input, stream);
Expand Down
64 changes: 42 additions & 22 deletions src/main/cpp/src/row_conversion.cu
View file
Original file line number Diff line number Diff line change
@@ -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.
Expand Down Expand Up @@ -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<size_type>(d_offsets_iterators.size());
Expand Down Expand Up @@ -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)};
}

Expand Down Expand Up @@ -1749,8 +1752,10 @@ std::vector<std::unique_ptr<column>> 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) {
Expand All @@ -1763,8 +1768,10 @@ std::vector<std::unique_ptr<column>> convert_to_rows(
thrust::make_transform_iterator(tbl.begin(), [](auto const &c) { return c.null_mask(); });
std::vector<bitmask_type const *> 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;
Expand Down Expand Up @@ -1810,7 +1817,8 @@ std::vector<std::unique_ptr<column>> 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();

Expand Down Expand Up @@ -1846,9 +1854,10 @@ std::vector<std::unique_ptr<column>> convert_to_rows(
std::vector<int8_t const *> 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];
Expand Down Expand Up @@ -2075,8 +2084,10 @@ std::unique_ptr<table> 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<std::unique_ptr<column>> output_columns;
Expand Down Expand Up @@ -2117,16 +2128,20 @@ std::unique_ptr<table> 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<detail::row_batch> row_batches;
row_batches.push_back(
{detail::row_batch{child.size(), num_rows, device_uvector<size_type>(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;
Expand Down Expand Up @@ -2163,7 +2178,8 @@ std::unique_ptr<table> 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());

Expand Down Expand Up @@ -2220,8 +2236,10 @@ std::unique_ptr<table> 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),
Expand Down Expand Up @@ -2273,8 +2291,10 @@ std::unique_ptr<table> 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<std::unique_ptr<column>> output_columns;
Expand Down
2 changes: 1 addition & 1 deletion thirdparty/cudf
Submodule cudf updated 107 files

0 comments on commit fea4a00

Please sign in to comment.