diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 98f6ed40502..56248f99f78 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -320,7 +320,8 @@ add_library(cudf src/strings/case.cu src/strings/char_types/char_cases.cu src/strings/char_types/char_types.cu - src/strings/combine.cu + src/strings/combine/concatenate.cu + src/strings/combine/join.cu src/strings/contains.cu src/strings/convert/convert_booleans.cu src/strings/convert/convert_datetime.cu diff --git a/cpp/src/strings/combine.cu b/cpp/src/strings/combine.cu deleted file mode 100644 index ebc31177b92..00000000000 --- a/cpp/src/strings/combine.cu +++ /dev/null @@ -1,476 +0,0 @@ -/* - * Copyright (c) 2019-2021, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include -#include - -#include -#include -#include - -#include - -namespace cudf { -namespace strings { -namespace detail { - -std::unique_ptr concatenate(table_view const& strings_columns, - string_scalar const& separator, - string_scalar const& narep, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - auto const num_columns = strings_columns.num_columns(); - CUDF_EXPECTS(num_columns > 0, "At least one column must be specified"); - // check all columns are of type string - CUDF_EXPECTS(std::all_of(strings_columns.begin(), - strings_columns.end(), - [](auto c) { return c.type().id() == type_id::STRING; }), - "All columns must be of type string"); - if (num_columns == 1) // single strings column returns a copy - return std::make_unique(*(strings_columns.begin()), stream, mr); - auto const strings_count = strings_columns.num_rows(); - if (strings_count == 0) // empty begets empty - return detail::make_empty_strings_column(stream, mr); - - CUDF_EXPECTS(separator.is_valid(), "Parameter separator must be a valid string_scalar"); - string_view d_separator(separator.data(), separator.size()); - auto d_narep = get_scalar_device_view(const_cast(narep)); - - // Create device views from the strings columns. - auto table = table_device_view::create(strings_columns, stream); - auto d_table = *table; - - // create resulting null mask - auto valid_mask = cudf::detail::valid_if( - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(strings_count), - [d_table, d_narep] __device__(size_type idx) { - bool null_element = thrust::any_of( - thrust::seq, d_table.begin(), d_table.end(), [idx](auto col) { return col.is_null(idx); }); - return (!null_element || d_narep.is_valid()); - }, - stream, - mr); - auto& null_mask = valid_mask.first; - auto const null_count = valid_mask.second; - - // build offsets column by computing sizes of each string in the output - auto offsets_transformer = [d_table, d_separator, d_narep] __device__(size_type row_idx) { - // for this row (idx), iterate over each column and add up the bytes - bool const null_element = - thrust::any_of(thrust::seq, d_table.begin(), d_table.end(), [row_idx](auto const& d_column) { - return d_column.is_null(row_idx); - }); - if (null_element && !d_narep.is_valid()) return 0; - size_type const bytes = thrust::transform_reduce( - thrust::seq, - d_table.begin(), - d_table.end(), - [row_idx, d_separator, d_narep] __device__(column_device_view const& d_column) { - return d_separator.size_bytes() + (d_column.is_null(row_idx) - ? d_narep.size() - : d_column.element(row_idx).size_bytes()); - }, - 0, - thrust::plus()); - // separator goes only in between elements - return bytes == 0 ? 0 : (bytes - d_separator.size_bytes()); // remove the last separator - }; - auto offsets_transformer_itr = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), offsets_transformer); - auto offsets_column = detail::make_offsets_child_column( - offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr); - auto d_results_offsets = offsets_column->view().data(); - - // create the chars column - auto const bytes = - cudf::detail::get_value(offsets_column->view(), strings_count, stream); - auto chars_column = strings::detail::create_chars_child_column(strings_count, bytes, stream, mr); - // fill the chars column - auto d_results_chars = chars_column->mutable_view().data(); - thrust::for_each_n( - rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - strings_count, - [d_table, num_columns, d_separator, d_narep, d_results_offsets, d_results_chars] __device__( - size_type idx) { - bool const null_element = thrust::any_of( - thrust::seq, d_table.begin(), d_table.end(), [idx](column_device_view const& col) { - return col.is_null(idx); - }); - if (null_element && !d_narep.is_valid()) - return; // do not write to buffer at all if any column element for this row is null - char* d_buffer = d_results_chars + d_results_offsets[idx]; - // write out each column's entry for this row - for (size_type col_idx = 0; col_idx < num_columns; ++col_idx) { - auto const d_column = d_table.column(col_idx); - string_view const d_str = - d_column.is_null(idx) ? d_narep.value() : d_column.element(idx); - d_buffer = detail::copy_string(d_buffer, d_str); - // separator goes only in between elements - if (col_idx + 1 < num_columns) d_buffer = detail::copy_string(d_buffer, d_separator); - } - }); - - return make_strings_column(strings_count, - std::move(offsets_column), - std::move(chars_column), - null_count, - std::move(null_mask), - stream, - mr); -} - -std::unique_ptr join_strings(strings_column_view const& strings, - string_scalar const& separator, - string_scalar const& narep, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - auto strings_count = strings.size(); - if (strings_count == 0) return detail::make_empty_strings_column(stream, mr); - - CUDF_EXPECTS(separator.is_valid(), "Parameter separator must be a valid string_scalar"); - - string_view d_separator(separator.data(), separator.size()); - auto d_narep = get_scalar_device_view(const_cast(narep)); - - auto strings_column = column_device_view::create(strings.parent(), stream); - auto d_strings = *strings_column; - - // create an offsets array for building the output memory layout - rmm::device_uvector output_offsets(strings_count + 1, stream); - auto d_output_offsets = output_offsets.data(); - // using inclusive-scan to compute last entry which is the total size - thrust::transform_inclusive_scan( - rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(strings_count), - d_output_offsets + 1, - [d_strings, d_separator, d_narep] __device__(size_type idx) { - size_type bytes = 0; - if (d_strings.is_null(idx)) { - if (!d_narep.is_valid()) return 0; // skip nulls - bytes += d_narep.size(); - } else - bytes += d_strings.element(idx).size_bytes(); - if ((idx + 1) < d_strings.size()) bytes += d_separator.size_bytes(); - return bytes; - }, - thrust::plus()); - size_type const zero = 0; - output_offsets.set_element_async(0, zero, stream); - // total size is the last entry - // Note this call does a synchronize on the stream and thereby also protects the - // set_element_async parameter from going out of scope before it is used. - size_type const bytes = output_offsets.back_element(stream); - - // build offsets column (only 1 string so 2 offset entries) - auto offsets_column = - make_numeric_column(data_type{type_id::INT32}, 2, mask_state::UNALLOCATED, stream, mr); - auto offsets_view = offsets_column->mutable_view(); - // set the first entry to 0 and the last entry to bytes - int32_t new_offsets[] = {0, static_cast(bytes)}; - CUDA_TRY(cudaMemcpyAsync(offsets_view.data(), - new_offsets, - sizeof(new_offsets), - cudaMemcpyHostToDevice, - stream.value())); - - // build null mask - // only one entry so it is either all valid or all null - size_type null_count = 0; - rmm::device_buffer null_mask{0, stream, mr}; // init to null null-mask - if (strings.null_count() == strings_count && !narep.is_valid()) { - null_mask = cudf::detail::create_null_mask(1, cudf::mask_state::ALL_NULL, stream, mr); - null_count = 1; - } - auto chars_column = detail::create_chars_child_column(strings_count, bytes, stream, mr); - auto chars_view = chars_column->mutable_view(); - auto d_chars = chars_view.data(); - thrust::for_each_n( - rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - strings_count, - [d_strings, d_separator, d_narep, d_output_offsets, d_chars] __device__(size_type idx) { - size_type offset = d_output_offsets[idx]; - char* d_buffer = d_chars + offset; - if (d_strings.is_null(idx)) { - if (!d_narep.is_valid()) - return; // do not write to buffer if element is null (including separator) - d_buffer = detail::copy_string(d_buffer, d_narep.value()); - } else { - string_view d_str = d_strings.element(idx); - d_buffer = detail::copy_string(d_buffer, d_str); - } - if ((idx + 1) < d_strings.size()) d_buffer = detail::copy_string(d_buffer, d_separator); - }); - - return make_strings_column(1, - std::move(offsets_column), - std::move(chars_column), - null_count, - std::move(null_mask), - stream, - mr); -} - -std::unique_ptr concatenate(table_view const& strings_columns, - strings_column_view const& separators, - string_scalar const& separator_narep, - string_scalar const& col_narep, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - auto const num_columns = strings_columns.num_columns(); - CUDF_EXPECTS(num_columns > 0, "At least one column must be specified"); - // Check if all columns are of type string - CUDF_EXPECTS(std::all_of(strings_columns.begin(), - strings_columns.end(), - [](auto c) { return c.type().id() == type_id::STRING; }), - "All columns must be of type string"); - - auto const strings_count = strings_columns.num_rows(); - CUDF_EXPECTS(strings_count == separators.size(), - "Separators column should be the same size as the strings columns"); - if (strings_count == 0) // Empty begets empty - return detail::make_empty_strings_column(stream, mr); - - // Invalid output column strings - null rows - string_view const invalid_str{nullptr, 0}; - auto const separator_rep = get_scalar_device_view(const_cast(separator_narep)); - auto const col_rep = get_scalar_device_view(const_cast(col_narep)); - auto const separator_col_view_ptr = column_device_view::create(separators.parent(), stream); - auto const separator_col_view = *separator_col_view_ptr; - - if (num_columns == 1) { - // Shallow copy of the resultant strings - rmm::device_uvector out_col_strings(strings_count, stream); - - // Device view of the only column in the table view - auto const col0_ptr = column_device_view::create(strings_columns.column(0), stream); - auto const col0 = *col0_ptr; - - // Execute it on every element - thrust::transform( - rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(strings_count), - out_col_strings.begin(), - // Output depends on the separator - [col0, invalid_str, separator_col_view, separator_rep, col_rep] __device__(auto ridx) { - if (!separator_col_view.is_valid(ridx) && !separator_rep.is_valid()) return invalid_str; - if (col0.is_valid(ridx)) { - auto sv = col0.element(ridx); - return sv.empty() ? string_view{} : sv; - } else if (col_rep.is_valid()) { - auto cv = col_rep.value(); - return cv.empty() ? string_view{} : cv; - } else - return invalid_str; - }); - - return make_strings_column(out_col_strings, invalid_str, stream, mr); - } - - // Create device views from the strings columns. - auto table = table_device_view::create(strings_columns, stream); - auto d_table = *table; - - // Create resulting null mask - auto valid_mask = cudf::detail::valid_if( - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(strings_count), - [d_table, separator_col_view, separator_rep, col_rep] __device__(size_type ridx) { - if (!separator_col_view.is_valid(ridx) && !separator_rep.is_valid()) return false; - bool all_nulls = - thrust::all_of(thrust::seq, d_table.begin(), d_table.end(), [ridx](auto const& col) { - return col.is_null(ridx); - }); - return all_nulls ? col_rep.is_valid() : true; - }, - stream, - mr); - - auto null_count = valid_mask.second; - - // Build offsets column by computing sizes of each string in the output - auto offsets_transformer = [d_table, separator_col_view, separator_rep, col_rep] __device__( - size_type ridx) { - // If the separator value for the row is null and if there aren't global separator - // replacements, this row does not have any value - null row - if (!separator_col_view.is_valid(ridx) && !separator_rep.is_valid()) return 0; - - // For this row (idx), iterate over each column and add up the bytes - bool const all_nulls = - thrust::all_of(thrust::seq, d_table.begin(), d_table.end(), [ridx](auto const& d_column) { - return d_column.is_null(ridx); - }); - // If all column values are null and there isn't a global column replacement value, this row - // is a null row - if (all_nulls && !col_rep.is_valid()) return 0; - - // There is at least one non-null column value (it can still be empty though) - auto const separator_str = separator_col_view.is_valid(ridx) - ? separator_col_view.element(ridx) - : separator_rep.value(); - - size_type const bytes = thrust::transform_reduce( - thrust::seq, - d_table.begin(), - d_table.end(), - [ridx, separator_str, col_rep] __device__(column_device_view const& d_column) { - // If column is null and there isn't a valid column replacement, this isn't used in - // final string concatenate - if (d_column.is_null(ridx) && !col_rep.is_valid()) return 0; - return separator_str.size_bytes() + (d_column.is_null(ridx) - ? col_rep.size() - : d_column.element(ridx).size_bytes()); - }, - 0, - thrust::plus()); - - // Null/empty separator and columns doesn't produce a non-empty string - if (bytes == 0) assert(separator_str.size_bytes() == 0); - - // Separator goes only in between elements - return static_cast(bytes - separator_str.size_bytes()); - }; - auto offsets_transformer_itr = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), offsets_transformer); - auto offsets_column = detail::make_offsets_child_column( - offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr); - auto d_results_offsets = offsets_column->view().data(); - - // Create the chars column - auto const bytes = - cudf::detail::get_value(offsets_column->view(), strings_count, stream); - auto chars_column = strings::detail::create_chars_child_column(strings_count, bytes, stream, mr); - - // Fill the chars column - auto d_results_chars = chars_column->mutable_view().data(); - thrust::for_each_n(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - strings_count, - [d_table, - num_columns, - d_results_offsets, - d_results_chars, - separator_col_view, - separator_rep, - col_rep] __device__(size_type ridx) { - // If the separator for this row is null and if there isn't a valid separator - // to replace, do not write anything for this row - if (!separator_col_view.is_valid(ridx) && !separator_rep.is_valid()) return; - - bool const all_nulls = thrust::all_of( - thrust::seq, d_table.begin(), d_table.end(), [ridx](auto const& col) { - return col.is_null(ridx); - }); - - // If all column values are null and there isn't a valid column replacement, - // skip this row - if (all_nulls && !col_rep.is_valid()) return; - - char* d_buffer = d_results_chars + d_results_offsets[ridx]; - bool colval_written = false; - - // There is at least one non-null column value (it can still be empty though) - auto const separator_str = separator_col_view.is_valid(ridx) - ? separator_col_view.element(ridx) - : separator_rep.value(); - - // Write out each column's entry for this row - for (size_type col_idx = 0; col_idx < num_columns; ++col_idx) { - auto const d_column = d_table.column(col_idx); - // If the row is null and if there is no replacement, skip it - if (d_column.is_null(ridx) && !col_rep.is_valid()) continue; - - // Separator goes only in between elements - if (colval_written) - d_buffer = detail::copy_string(d_buffer, separator_str); - - string_view const d_str = d_column.is_null(ridx) - ? col_rep.value() - : d_column.element(ridx); - d_buffer = detail::copy_string(d_buffer, d_str); - colval_written = true; - } - }); - - return make_strings_column(strings_count, - std::move(offsets_column), - std::move(chars_column), - null_count, - (null_count) ? std::move(valid_mask.first) : rmm::device_buffer{}, - stream, - mr); -} - -} // namespace detail - -// APIs - -std::unique_ptr concatenate(table_view const& strings_columns, - string_scalar const& separator, - string_scalar const& narep, - rmm::mr::device_memory_resource* mr) -{ - CUDF_FUNC_RANGE(); - return detail::concatenate(strings_columns, separator, narep, rmm::cuda_stream_default, mr); -} - -std::unique_ptr join_strings(strings_column_view const& strings, - string_scalar const& separator, - string_scalar const& narep, - rmm::mr::device_memory_resource* mr) -{ - CUDF_FUNC_RANGE(); - return detail::join_strings(strings, separator, narep, rmm::cuda_stream_default, mr); -} - -std::unique_ptr concatenate(table_view const& strings_columns, - strings_column_view const& separators, - string_scalar const& separator_narep, - string_scalar const& col_narep, - rmm::mr::device_memory_resource* mr) -{ - CUDF_FUNC_RANGE(); - return detail::concatenate( - strings_columns, separators, separator_narep, col_narep, rmm::cuda_stream_default, mr); -} - -} // namespace strings -} // namespace cudf diff --git a/cpp/src/strings/combine/concatenate.cu b/cpp/src/strings/combine/concatenate.cu new file mode 100644 index 00000000000..40ce39192d8 --- /dev/null +++ b/cpp/src/strings/combine/concatenate.cu @@ -0,0 +1,290 @@ +/* + * Copyright (c) 2019-2021, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include + +#include + +#include + +namespace cudf { +namespace strings { +namespace detail { +namespace { + +/** + * @brief Concatenate strings functor + * + * This will concatenate the strings from each row of the given table + * and apply the separator. The null-replacement string `d_narep` is + * used in place of any string in a row that contains a null entry. + */ +struct concat_strings_fn { + table_device_view const d_table; + string_view const d_separator; + string_scalar_device_view const d_narep; + offset_type* d_offsets{}; + char* d_chars{}; + + __device__ void operator()(size_type idx) + { + bool const null_element = + thrust::any_of(thrust::seq, d_table.begin(), d_table.end(), [idx](auto const& col) { + return col.is_null(idx); + }); + // handle a null row + if (null_element && !d_narep.is_valid()) { + if (!d_chars) d_offsets[idx] = 0; + return; + } + + char* d_buffer = d_chars ? d_chars + d_offsets[idx] : nullptr; + size_type bytes = 0; + for (auto itr = d_table.begin(); itr < d_table.end(); ++itr) { + auto const d_column = *itr; + auto const d_str = + d_column.is_null(idx) ? d_narep.value() : d_column.element(idx); + if (d_buffer) d_buffer = detail::copy_string(d_buffer, d_str); + bytes += d_str.size_bytes(); + // separator goes only in between elements + if (itr + 1 < d_table.end()) { + if (d_buffer) d_buffer = detail::copy_string(d_buffer, d_separator); + bytes += d_separator.size_bytes(); + } + } + if (!d_chars) d_offsets[idx] = bytes; + } +}; + +} // namespace + +std::unique_ptr concatenate(table_view const& strings_columns, + string_scalar const& separator, + string_scalar const& narep, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + auto const num_columns = strings_columns.num_columns(); + CUDF_EXPECTS(num_columns > 0, "At least one column must be specified"); + // check all columns are of type string + CUDF_EXPECTS(std::all_of(strings_columns.begin(), + strings_columns.end(), + [](auto c) { return c.type().id() == type_id::STRING; }), + "All columns must be of type string"); + if (num_columns == 1) // single strings column returns a copy + return std::make_unique(*(strings_columns.begin()), stream, mr); + auto const strings_count = strings_columns.num_rows(); + if (strings_count == 0) // empty begets empty + return detail::make_empty_strings_column(stream, mr); + + CUDF_EXPECTS(separator.is_valid(), "Parameter separator must be a valid string_scalar"); + string_view d_separator(separator.data(), separator.size()); + auto d_narep = get_scalar_device_view(const_cast(narep)); + + // Create device views from the strings columns. + auto d_table = table_device_view::create(strings_columns, stream); + concat_strings_fn fn{*d_table, d_separator, d_narep}; + auto children = make_strings_children(fn, strings_count, stream, mr); + + // create resulting null mask + auto [null_mask, null_count] = cudf::detail::valid_if( + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(strings_count), + [d_table = *d_table, d_narep] __device__(size_type idx) { + bool null_element = thrust::any_of( + thrust::seq, d_table.begin(), d_table.end(), [idx](auto col) { return col.is_null(idx); }); + return (!null_element || d_narep.is_valid()); + }, + stream, + mr); + + return make_strings_column(strings_count, + std::move(children.first), + std::move(children.second), + null_count, + std::move(null_mask), + stream, + mr); +} + +namespace { + +/** + * @brief Concatenate strings functor using multiple separators. + * + * A unique separator is provided for each row along with a string to use + * when a separator row is null `d_separator_narep`. The `d_narep` is + * used in place of a null entry in the strings columns. + */ +struct multi_separator_concat_fn { + table_device_view const d_table; + column_device_view const d_separators; + string_scalar_device_view const d_separator_narep; + string_scalar_device_view const d_narep; + offset_type* d_offsets{}; + char* d_chars{}; + + __device__ void operator()(size_type idx) + { + bool const all_nulls = + thrust::all_of(thrust::seq, d_table.begin(), d_table.end(), [idx](auto const& col) { + return col.is_null(idx); + }); + + if ((d_separators.is_null(idx) && !d_separator_narep.is_valid()) || + (all_nulls && !d_narep.is_valid())) { + if (!d_chars) d_offsets[idx] = 0; + return; + } + + // point to output location + char* d_buffer = d_chars ? d_chars + d_offsets[idx] : nullptr; + offset_type bytes = 0; + + // there is at least one non-null column value + auto const d_separator = d_separators.is_valid(idx) ? d_separators.element(idx) + : d_separator_narep.value(); + auto const d_null_rep = d_narep.is_valid() ? d_narep.value() : string_view{}; + + // write output entry for this row + bool colval_written = false; // state variable for writing separators + for (auto const d_column : d_table) { + // if the row is null and if there is no replacement, skip it + if (d_column.is_null(idx) && !d_narep.is_valid()) continue; + + // separator in this row is written only after the first output + if (colval_written) { + if (d_buffer) d_buffer = detail::copy_string(d_buffer, d_separator); + bytes += d_separator.size_bytes(); + } + + // write out column's row data (or narep if the row is null) + string_view const d_str = + d_column.is_null(idx) ? d_null_rep : d_column.element(idx); + if (d_buffer) d_buffer = detail::copy_string(d_buffer, d_str); + bytes += d_str.size_bytes(); + + // column's string or narep could by empty so we need this flag + // to know we got this far even if no actual bytes were copied + colval_written = true; // use the separator before the next column + } + + if (!d_chars) d_offsets[idx] = bytes; + } +}; +} // namespace + +std::unique_ptr concatenate(table_view const& strings_columns, + strings_column_view const& separators, + string_scalar const& separator_narep, + string_scalar const& col_narep, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + auto const num_columns = strings_columns.num_columns(); + CUDF_EXPECTS(num_columns > 0, "At least one column must be specified"); + // Check if all columns are of type string + CUDF_EXPECTS(std::all_of(strings_columns.begin(), + strings_columns.end(), + [](auto c) { return c.type().id() == type_id::STRING; }), + "All columns must be of type string"); + + auto const strings_count = strings_columns.num_rows(); + CUDF_EXPECTS(strings_count == separators.size(), + "Separators column should be the same size as the strings columns"); + if (strings_count == 0) // Empty begets empty + return detail::make_empty_strings_column(stream, mr); + + // Invalid output column strings - null rows + string_view const invalid_str{nullptr, 0}; + auto const separator_rep = get_scalar_device_view(const_cast(separator_narep)); + auto const col_rep = get_scalar_device_view(const_cast(col_narep)); + auto const separator_col_view_ptr = column_device_view::create(separators.parent(), stream); + auto const separator_col_view = *separator_col_view_ptr; + + // Create device views from the strings columns. + auto d_table = table_device_view::create(strings_columns, stream); + + multi_separator_concat_fn mscf{*d_table, separator_col_view, separator_rep, col_rep}; + auto children = make_strings_children(mscf, strings_count, stream, mr); + + // Create resulting null mask + auto [null_mask, null_count] = cudf::detail::valid_if( + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(strings_count), + [d_table = *d_table, separator_col_view, separator_rep, col_rep] __device__(size_type ridx) { + if (!separator_col_view.is_valid(ridx) && !separator_rep.is_valid()) return false; + bool all_nulls = + thrust::all_of(thrust::seq, d_table.begin(), d_table.end(), [ridx](auto const& col) { + return col.is_null(ridx); + }); + return all_nulls ? col_rep.is_valid() : true; + }, + stream, + mr); + + return make_strings_column(strings_count, + std::move(children.first), + std::move(children.second), + null_count, + std::move(null_mask), + stream, + mr); +} + +} // namespace detail + +// APIs + +std::unique_ptr concatenate(table_view const& strings_columns, + string_scalar const& separator, + string_scalar const& narep, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::concatenate(strings_columns, separator, narep, rmm::cuda_stream_default, mr); +} + +std::unique_ptr concatenate(table_view const& strings_columns, + strings_column_view const& separators, + string_scalar const& separator_narep, + string_scalar const& col_narep, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::concatenate( + strings_columns, separators, separator_narep, col_narep, rmm::cuda_stream_default, mr); +} + +} // namespace strings +} // namespace cudf diff --git a/cpp/src/strings/combine/join.cu b/cpp/src/strings/combine/join.cu new file mode 100644 index 00000000000..8faec80ee71 --- /dev/null +++ b/cpp/src/strings/combine/join.cu @@ -0,0 +1,147 @@ +/* + * Copyright (c) 2019-2021, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include + +#include + +namespace cudf { +namespace strings { +namespace detail { + +std::unique_ptr join_strings(strings_column_view const& strings, + string_scalar const& separator, + string_scalar const& narep, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + auto strings_count = strings.size(); + if (strings_count == 0) return detail::make_empty_strings_column(stream, mr); + + CUDF_EXPECTS(separator.is_valid(), "Parameter separator must be a valid string_scalar"); + + string_view d_separator(separator.data(), separator.size()); + auto d_narep = get_scalar_device_view(const_cast(narep)); + + auto strings_column = column_device_view::create(strings.parent(), stream); + auto d_strings = *strings_column; + + // create an offsets array for building the output memory layout + rmm::device_uvector output_offsets(strings_count + 1, stream); + auto d_output_offsets = output_offsets.data(); + // using inclusive-scan to compute last entry which is the total size + thrust::transform_inclusive_scan( + rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(strings_count), + d_output_offsets + 1, + [d_strings, d_separator, d_narep] __device__(size_type idx) { + size_type bytes = 0; + if (d_strings.is_null(idx)) { + if (!d_narep.is_valid()) return 0; // skip nulls + bytes += d_narep.size(); + } else + bytes += d_strings.element(idx).size_bytes(); + if ((idx + 1) < d_strings.size()) bytes += d_separator.size_bytes(); + return bytes; + }, + thrust::plus()); + size_type const zero = 0; + output_offsets.set_element_async(0, zero, stream); + // total size is the last entry + // Note this call does a synchronize on the stream and thereby also protects the + // set_element_async parameter from going out of scope before it is used. + size_type const bytes = output_offsets.back_element(stream); + + // build offsets column (only 1 string so 2 offset entries) + auto offsets_column = + make_numeric_column(data_type{type_id::INT32}, 2, mask_state::UNALLOCATED, stream, mr); + auto offsets_view = offsets_column->mutable_view(); + // set the first entry to 0 and the last entry to bytes + int32_t new_offsets[] = {0, static_cast(bytes)}; + CUDA_TRY(cudaMemcpyAsync(offsets_view.data(), + new_offsets, + sizeof(new_offsets), + cudaMemcpyHostToDevice, + stream.value())); + + // build null mask + // only one entry so it is either all valid or all null + auto const null_count = + static_cast(strings.null_count() == strings_count && !narep.is_valid()); + auto null_mask = null_count + ? cudf::detail::create_null_mask(1, cudf::mask_state::ALL_NULL, stream, mr) + : rmm::device_buffer{0, stream, mr}; + auto chars_column = detail::create_chars_child_column(strings_count, bytes, stream, mr); + auto d_chars = chars_column->mutable_view().data(); + thrust::for_each_n( + rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + strings_count, + [d_strings, d_separator, d_narep, d_output_offsets, d_chars] __device__(size_type idx) { + size_type offset = d_output_offsets[idx]; + char* d_buffer = d_chars + offset; + if (d_strings.is_null(idx)) { + if (!d_narep.is_valid()) + return; // do not write to buffer if element is null (including separator) + d_buffer = detail::copy_string(d_buffer, d_narep.value()); + } else { + string_view d_str = d_strings.element(idx); + d_buffer = detail::copy_string(d_buffer, d_str); + } + if ((idx + 1) < d_strings.size()) d_buffer = detail::copy_string(d_buffer, d_separator); + }); + + return make_strings_column(1, + std::move(offsets_column), + std::move(chars_column), + null_count, + std::move(null_mask), + stream, + mr); +} + +} // namespace detail + +// external API + +std::unique_ptr join_strings(strings_column_view const& strings, + string_scalar const& separator, + string_scalar const& narep, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::join_strings(strings, separator, narep, rmm::cuda_stream_default, mr); +} + +} // namespace strings +} // namespace cudf diff --git a/cpp/tests/strings/combine_tests.cpp b/cpp/tests/strings/combine_tests.cpp index cfeca2bba29..2ca0562064d 100644 --- a/cpp/tests/strings/combine_tests.cpp +++ b/cpp/tests/strings/combine_tests.cpp @@ -265,7 +265,7 @@ TEST_F(StringsConcatenateWithColSeparatorTest, auto results = cudf::strings::concatenate( cudf::table_view{{col0}}, cudf::strings_column_view(sep_col), sep_rep, col_rep); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, exp_results, true); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, exp_results, true); } TEST_F(StringsConcatenateWithColSeparatorTest, SingleColumnStringMixNoReplacements) @@ -343,7 +343,7 @@ TEST_F(StringsConcatenateWithColSeparatorTest, SingleColumnStringMixSeparatorAnd auto results = cudf::strings::concatenate( cudf::table_view{{col0}}, cudf::strings_column_view(sep_col), sep_rep, col_rep); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, exp_results, true); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, exp_results, true); } TEST_F(StringsConcatenateWithColSeparatorTest, MultiColumnEmptyAndNullStringsNoReplacements) @@ -483,7 +483,7 @@ TEST_F(StringsConcatenateWithColSeparatorTest, MultiColumnStringMixSeparatorAndC auto results = cudf::strings::concatenate( cudf::table_view{{col0, col1}}, cudf::strings_column_view(sep_col), sep_rep, col_rep); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, exp_results, true); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, exp_results, true); } TEST_F(StringsConcatenateWithColSeparatorTest, MultiColumnNonNullableStrings) @@ -499,5 +499,5 @@ TEST_F(StringsConcatenateWithColSeparatorTest, MultiColumnNonNullableStrings) auto results = cudf::strings::concatenate(cudf::table_view{{col0, col1}}, cudf::strings_column_view(sep_col)); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, exp_results, true); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, exp_results, true); }