diff --git a/conda/recipes/libcudf/meta.yaml b/conda/recipes/libcudf/meta.yaml index 885a22870bb..a46712def28 100644 --- a/conda/recipes/libcudf/meta.yaml +++ b/conda/recipes/libcudf/meta.yaml @@ -96,6 +96,7 @@ test: - test -f $PREFIX/include/cudf/detail/utilities/alignment.hpp - test -f $PREFIX/include/cudf/detail/utilities/integer_utils.hpp - test -f $PREFIX/include/cudf/detail/utilities/int_fastdiv.h + - test -f $PREFIX/include/cudf/detail/utilities/vector_factories.hpp - test -f $PREFIX/include/cudf/dictionary/detail/concatenate.hpp - test -f $PREFIX/include/cudf/dictionary/detail/encode.hpp - test -f $PREFIX/include/cudf/dictionary/detail/merge.hpp diff --git a/cpp/benchmarks/common/generate_benchmark_input.cpp b/cpp/benchmarks/common/generate_benchmark_input.cpp index f20d2cab725..a66416ad40b 100644 --- a/cpp/benchmarks/common/generate_benchmark_input.cpp +++ b/cpp/benchmarks/common/generate_benchmark_input.cpp @@ -26,6 +26,7 @@ #include #include +#include #include #include @@ -411,7 +412,11 @@ std::unique_ptr create_random_column(data_profi row += std::max(run_len - 1, 0); } } - return cudf::make_strings_column(out_col.chars, out_col.offsets, out_col.null_mask); + + rmm::device_vector d_chars(out_col.chars); + rmm::device_vector d_offsets(out_col.offsets); + rmm::device_vector d_null_mask(out_col.null_mask); + return cudf::make_strings_column(d_chars, d_offsets, d_null_mask); } template <> diff --git a/cpp/benchmarks/copying/shift_benchmark.cu b/cpp/benchmarks/copying/shift_benchmark.cu index 291c0ef6777..42d8b58aca3 100644 --- a/cpp/benchmarks/copying/shift_benchmark.cu +++ b/cpp/benchmarks/copying/shift_benchmark.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -23,7 +23,6 @@ #include -#include #include #include #include diff --git a/cpp/include/cudf/column/column_factories.hpp b/cpp/include/cudf/column/column_factories.hpp index 7ccc5879f5f..31196824845 100644 --- a/cpp/include/cudf/column/column_factories.hpp +++ b/cpp/include/cudf/column/column_factories.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -17,6 +17,7 @@ #include #include +#include #include #include @@ -41,9 +42,8 @@ namespace cudf { std::unique_ptr make_empty_column(data_type type); /** - * @brief Construct column with sufficient uninitialized storage - * to hold `size` elements of the specified numeric `data_type` with an optional - * null mask. + * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the + * specified numeric `data_type` with an optional null mask. * * @note `null_count()` is determined by the requested null mask `state` * @@ -65,9 +65,8 @@ std::unique_ptr make_numeric_column( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Construct column with sufficient uninitialized storage - * to hold `size` elements of the specified numeric `data_type` with a - * null mask. + * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the + * specified numeric `data_type` with a null mask. * * @note null_count is optional and will be computed if not provided. * @@ -153,9 +152,8 @@ std::unique_ptr make_fixed_point_column( } /** - * @brief Construct column with sufficient uninitialized storage - * to hold `size` elements of the specified timestamp `data_type` with an - * optional null mask. + * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the + * specified timestamp `data_type` with an optional null mask. * * @note `null_count()` is determined by the requested null mask `state` * @@ -177,9 +175,8 @@ std::unique_ptr make_timestamp_column( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Construct column with sufficient uninitialized storage - * to hold `size` elements of the specified timestamp `data_type` with a - * null mask. + * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the + * specified timestamp `data_type` with a null mask. * * @note null_count is optional and will be computed if not provided. * @@ -211,9 +208,8 @@ std::unique_ptr make_timestamp_column( } /** - * @brief Construct column with sufficient uninitialized storage - * to hold `size` elements of the specified duration `data_type` with an - * optional null mask. + * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the + * specified duration `data_type` with an optional null mask. * * @note `null_count()` is determined by the requested null mask `state` * @@ -235,9 +231,8 @@ std::unique_ptr make_duration_column( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Construct column with sufficient uninitialized storage - * to hold `size` elements of the specified duration `data_type` with a - * null mask. + * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the + * specified duration `data_type` with a null mask. * * @note null_count is optional and will be computed if not provided. * @@ -269,9 +264,8 @@ std::unique_ptr make_duration_column( } /** - * @brief Construct column with sufficient uninitialized storage - * to hold `size` elements of the specified fixed width `data_type` with an optional - * null mask. + * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the + * specified fixed width `data_type` with an optional null mask. * * @note `null_count()` is determined by the requested null mask `state` * @@ -293,9 +287,8 @@ std::unique_ptr make_fixed_width_column( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Construct column with sufficient uninitialized storage - * to hold `size` elements of the specified fixed width `data_type` with a - * null mask. + * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the + * specified fixed width `data_type` with a null mask. * * @note null_count is optional and will be computed if not provided. * @@ -330,7 +323,8 @@ std::unique_ptr make_fixed_width_column( } /** - * @brief Construct STRING type column given a device vector of pointer/size pairs. + * @brief Construct a STRING type column given a device span of pointer/size pairs. + * * The total number of char bytes must not exceed the maximum size of size_type. * The string characters are expected to be UTF-8 encoded sequence of char * bytes. Use the strings_column_view class to perform strings operations on @@ -344,20 +338,20 @@ std::unique_ptr make_fixed_width_column( * * @throws std::bad_alloc if device memory allocation fails * - * @param[in] strings The vector of pointer/size pairs. - * Each pointer must be a device memory address or `nullptr` - * (indicating a null string). The size must be the number of bytes. + * @param[in] strings The device span of pointer/size pairs. Each pointer must be a device memory + address or `nullptr` (indicating a null string). The size must be the number of bytes. * @param[in] stream CUDA stream used for device memory operations and kernel launches. * @param[in] mr Device memory resource used for allocation of the column's `null_mask` and children * columns' device memory. */ std::unique_ptr make_strings_column( - const rmm::device_vector>& strings, + cudf::device_span const> strings, rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Construct STRING type column given a device vector of string_view. + * @brief Construct a STRING type column given a device span of string_view. + * * The total number of char bytes must not exceed the maximum size of size_type. * The string characters are expected to be UTF-8 encoded sequence of char * bytes. Use the strings_column_view class to perform strings operations on @@ -372,10 +366,8 @@ std::unique_ptr make_strings_column( * * @throws std::bad_alloc if device memory allocation fails * - * @param[in] string_views The vector of string_view. - * Each string_view must point to a device memory address or - * `null_placeholder` (indicating a null string). The size must be the number of - * bytes. + * @param[in] string_views The span of string_view. Each string_view must point to a device memory + address or `null_placeholder` (indicating a null string). The size must be the number of bytes. * @param[in] null_placeholder string_view indicating null string in given list of * string_views. * @param[in] stream CUDA stream used for device memory operations and kernel launches. @@ -383,107 +375,61 @@ std::unique_ptr make_strings_column( * columns' device memory. */ std::unique_ptr make_strings_column( - const rmm::device_vector& string_views, + cudf::device_span string_views, const string_view null_placeholder, rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Construct STRING type column given a device vector of chars - * encoded as UTF-8, a device vector of byte offsets identifying individual - * strings within the char vector, and an optional null bitmask. + * @brief Construct a STRING type column given a device span of chars encoded as UTF-8, a device + * span of byte offsets identifying individual strings within the char vector, and an optional + * null bitmask. * * `offsets.front()` must always be zero. * - * The total number of char bytes must not exceed the maximum size of size_type. - * Use the strings_column_view class to perform strings operations on this type - * of column. - * This function makes a deep copy of the strings, offsets, null_mask to create - * a new column. + * The total number of char bytes must not exceed the maximum size of size_type. Use the + * strings_column_view class to perform strings operations on this type of column. * - * @throws std::bad_alloc if device memory allocation fails - * - * @param[in] strings The vector of chars in device memory. - * This char vector is expected to be UTF-8 encoded characters. - * @param[in] offsets The vector of byte offsets in device memory. - * The number of elements is one more than the total number - * of strings so the `offsets.back()` is the total - * number of bytes in the strings array. - * `offsets.front()` must always be 0 to point to the beginning - * of `strings`. - * @param[in] null_mask Device vector containing the null element indicator bitmask. - * Arrow format for nulls is used for interpeting this bitmask. - * @param[in] null_count The number of null string entries. If equal to - * `UNKNOWN_NULL_COUNT`, the null count will be computed dynamically on the - * first invocation of `column::null_count()` - * @param[in] stream CUDA stream used for device memory operations and kernel launches. - * @param[in] mr Device memory resource used for allocation of the column's `null_mask` and children - * columns' device memory. - */ -std::unique_ptr make_strings_column( - const rmm::device_vector& strings, - const rmm::device_vector& offsets, - const rmm::device_vector& null_mask = {}, - size_type null_count = cudf::UNKNOWN_NULL_COUNT, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); - -/** - * @brief Construct STRING type column given a host vector of chars - * encoded as UTF-8, a host vector of byte offsets identifying individual - * strings within the char vector, and an optional null bitmask. - * - * `offsets.front()` must always be zero. - * - * The total number of char bytes must not exceed the maximum size of size_type. - * Use the strings_column_view class to perform strings operations on this type - * of column. - * This function makes a deep copy of the strings, offsets, null_mask to create - * a new column. + * This function makes a deep copy of the strings, offsets, null_mask to create a new column. * * @throws std::bad_alloc if device memory allocation fails * - * @param[in] strings The contiguous array of chars in host memory. - * This char array is expected to be UTF-8 encoded characters. - * @param[in] offsets The array of byte offsets in host memory. - * The number of elements is one more than the total number - * of strings so the `offsets.back()` is the total - * number of bytes in the strings array. - * `offsets.front()` must always be 0 to point to the beginning - * of `strings`. - * @param[in] null_mask Host vector containing the null element indicator bitmask. - * Arrow format for nulls is used for interpeting this bitmask. - * @param[in] null_count The number of null string entries. If equal to - * `UNKNOWN_NULL_COUNT`, the null count will be computed dynamically on the - * first invocation of `column::null_count()` + * @param[in] strings The device span of chars in device memory. This char vector is expected to be + * UTF-8 encoded characters. + * @param[in] offsets The device span of byte offsets in device memory. The number of elements is + * one more than the total number of strings so the `offsets.back()` is the total number of bytes + * in the strings array. `offsets.front()` must always be 0 to point to the beginning of `strings`. + * @param[in] null_mask Device span containing the null element indicator bitmask. Arrow format for + * nulls is used for interpeting this bitmask. + * @param[in] null_count The number of null string entries. If equal to `UNKNOWN_NULL_COUNT`, the + * null count will be computed dynamically on the first invocation of `column::null_count()` * @param[in] stream CUDA stream used for device memory operations and kernel launches. * @param[in] mr Device memory resource used for allocation of the column's `null_mask` and children * columns' device memory. */ std::unique_ptr make_strings_column( - const std::vector& strings, - const std::vector& offsets, - const std::vector& null_mask = {}, - size_type null_count = cudf::UNKNOWN_NULL_COUNT, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + cudf::device_span strings, + cudf::device_span offsets, + cudf::device_span null_mask = {}, + size_type null_count = cudf::UNKNOWN_NULL_COUNT, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Constructs a STRING type column given offsets column, chars columns, - * and null mask and null count. The columns and mask are moved into the - * resulting strings column. + * @brief Construct a STRING type column given offsets column, chars columns, and null mask and null + * count. + * + * The columns and mask are moved into the resulting strings column. * * @param[in] num_strings The number of strings the column represents. - * @param[in] offsets_column The column of offset values for this column. - * The number of elements is one more than the total number - * of strings so the offset[last] - offset[0] is the total - * number of bytes in the strings vector. - * @param[in] chars_column The column of char bytes for all the strings for this column. - * Individual strings are identified by the offsets and the - * nullmask. + * @param[in] offsets_column The column of offset values for this column. The number of elements is + * one more than the total number of strings so the `offset[last] - offset[0]` is the total number + * of bytes in the strings vector. + * @param[in] chars_column The column of char bytes for all the strings for this column. Individual + * strings are identified by the offsets and the nullmask. * @param[in] null_count The number of null string entries. - * @param[in] null_mask The bits specifying the null strings in device memory. - * Arrow format for nulls is used for interpeting this bitmask. + * @param[in] null_mask The bits specifying the null strings in device memory. Arrow format for + * nulls is used for interpeting this bitmask. * @param[in] stream CUDA stream used for device memory operations and kernel launches. * @param[in] mr Device memory resource used for allocation of the column's `null_mask` and children * columns' device memory. @@ -498,12 +444,11 @@ std::unique_ptr make_strings_column( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Constructs a LIST type column given offsets column, child column, - * and null mask and null count. + * @brief Construct a LIST type column given offsets column, child column, null mask and null + * count. * * The columns and mask are moved into the resulting lists column. * - * * List columns are structured similarly to strings columns. They contain * a set of offsets which represents the lengths of the lists in each row, and * a "child" column of data that is referenced by the offsets. Since lists @@ -563,7 +508,7 @@ std::unique_ptr make_lists_column( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Constructs a STRUCT column using specified child columns as members. + * @brief Construct a STRUCT column using specified child columns as members. * * Specified child/member columns and null_mask are adopted by resultant * struct column. @@ -593,8 +538,7 @@ std::unique_ptr make_structs_column( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Return a column with size elements that are all equal to the - * given scalar. + * @brief Construct a column with size elements that are all equal to the given scalar. * * The output column will have the same type as `s.type()` * The output column will contain all null rows if `s.invalid()==false` @@ -612,8 +556,7 @@ std::unique_ptr make_column_from_scalar( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Return a dictionary column with size elements that are all equal to the - * given scalar. + * @brief Construct a dictionary column with size elements that are all equal to the given scalar. * * The output column will have keys of type `s.type()` * The output column will be empty if `size==0`. diff --git a/cpp/include/cudf/detail/utilities/trie.cuh b/cpp/include/cudf/detail/utilities/trie.cuh index 77b184a4874..f2d429d5529 100644 --- a/cpp/include/cudf/detail/utilities/trie.cuh +++ b/cpp/include/cudf/detail/utilities/trie.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2020, NVIDIA CORPORATION. + * Copyright (c) 2018-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. @@ -30,7 +30,7 @@ #include #include -using cudf::detail::device_span; +using cudf::device_span; static constexpr char trie_terminating_character = '\n'; diff --git a/cpp/include/cudf/detail/utilities/vector_factories.hpp b/cpp/include/cudf/detail/utilities/vector_factories.hpp new file mode 100644 index 00000000000..030d2c331c5 --- /dev/null +++ b/cpp/include/cudf/detail/utilities/vector_factories.hpp @@ -0,0 +1,236 @@ +/* + * Copyright (c) 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. + */ + +/** + * @brief Convenience factories for creating device vectors from host spans + * @file vector_factories.hpp + */ + +#include + +#include +#include +#include + +namespace cudf { +namespace detail { + +/** + * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a + * `host_span` + * + * @note This function does not synchronize `stream`. + * + * @tparam T The type of the data to copy + * @param source_data The host_span of data to deep copy + * @param stream The stream on which to allocate memory and perform the copy + * @param mr The memory resource to use for allocating the returned device_uvector + * @return A device_uvector containing the copied data + */ +template +rmm::device_uvector make_device_uvector_async( + host_span source_data, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + rmm::device_uvector ret(source_data.size(), stream, mr); + CUDA_TRY(cudaMemcpyAsync(ret.data(), + source_data.data(), + source_data.size() * sizeof(T), + cudaMemcpyDefault, + stream.value())); + return ret; +} + +/** + * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a host + * container + * + * @note This function does not synchronize `stream`. + * + * @tparam Container The type of the container to copy from + * @tparam T The type of the data to copy + * @param c The input host container from which to copy + * @param stream The stream on which to allocate memory and perform the copy + * @param mr The memory resource to use for allocating the returned device_uvector + * @return A device_uvector containing the copied data + */ +template >::value>* = nullptr> +rmm::device_uvector make_device_uvector_async( + Container const& c, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + return make_device_uvector_async(host_span{c}, stream, mr); +} + +/** + * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a + * `device_span` + * + * @note This function does not synchronize `stream`. + * + * @tparam T The type of the data to copy + * @param source_data The device_span of data to deep copy + * @param stream The stream on which to allocate memory and perform the copy + * @param mr The memory resource to use for allocating the returned device_uvector + * @return A device_uvector containing the copied data + */ +template +rmm::device_uvector make_device_uvector_async( + device_span source_data, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + rmm::device_uvector ret(source_data.size(), stream, mr); + CUDA_TRY(cudaMemcpyAsync(ret.data(), + source_data.data(), + source_data.size() * sizeof(T), + cudaMemcpyDefault, + stream.value())); + return ret; +} + +/** + * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a device + * container + * + * @note This function does not synchronize `stream`. + * + * @tparam Container The type of the container to copy from + * @tparam T The type of the data to copy + * @param c The input device container from which to copy + * @param stream The stream on which to allocate memory and perform the copy + * @param mr The memory resource to use for allocating the returned device_uvector + * @return A device_uvector containing the copied data + */ +template < + typename Container, + std::enable_if_t< + std::is_convertible>::value>* = + nullptr> +rmm::device_uvector make_device_uvector_async( + Container const& c, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + return make_device_uvector_async( + device_span{c}, stream, mr); +} + +/** + * @brief Synchronously construct a `device_uvector` containing a deep copy of data from a + * `host_span` + * + * @note This function synchronizes `stream`. + * + * @tparam T The type of the data to copy + * @param source_data The host_span of data to deep copy + * @param stream The stream on which to allocate memory and perform the copy + * @param mr The memory resource to use for allocating the returned device_uvector + * @return A device_uvector containing the copied data + */ +template +rmm::device_uvector make_device_uvector_sync( + host_span source_data, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + auto ret = make_device_uvector_async(source_data, stream, mr); + stream.synchronize(); + return ret; +} + +/** + * @brief Synchronously construct a `device_uvector` containing a deep copy of data from a host + * container + * + * @note This function synchronizes `stream`. + * + * @tparam Container The type of the container to copy from + * @tparam T The type of the data to copy + * @param c The input host container from which to copy + * @param stream The stream on which to allocate memory and perform the copy + * @param mr The memory resource to use for allocating the returned device_uvector + * @return A device_uvector containing the copied data + */ +template >::value>* = nullptr> +rmm::device_uvector make_device_uvector_sync( + Container const& c, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + return make_device_uvector_sync(host_span{c}, stream, mr); +} + +/** + * @brief Synchronously construct a `device_uvector` containing a deep copy of data from a + * `device_span` + * + * @note This function synchronizes `stream`. + * + * @tparam T The type of the data to copy + * @param source_data The device_span of data to deep copy + * @param stream The stream on which to allocate memory and perform the copy + * @param mr The memory resource to use for allocating the returned device_uvector + * @return A device_uvector containing the copied data + */ +template +rmm::device_uvector make_device_uvector_sync( + device_span source_data, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + auto ret = make_device_uvector_async(source_data, stream, mr); + stream.synchronize(); + return ret; +} + +/** + * @brief Synchronously construct a `device_uvector` containing a deep copy of data from a device + * container + * + * @note This function synchronizes `stream`. + * + * @tparam Container The type of the container to copy from + * @tparam T The type of the data to copy + * @param c The input device container from which to copy + * @param stream The stream on which to allocate memory and perform the copy + * @param mr The memory resource to use for allocating the returned device_uvector + * @return A device_uvector containing the copied data + */ +template < + typename Container, + std::enable_if_t< + std::is_convertible>::value>* = + nullptr> +rmm::device_uvector make_device_uvector_sync( + Container const& c, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + return make_device_uvector_sync(device_span{c}, stream, mr); +} + +} // namespace detail + +} // namespace cudf diff --git a/cpp/include/cudf/strings/detail/gather.cuh b/cpp/include/cudf/strings/detail/gather.cuh index cbc8ea244d2..28da8ef4324 100644 --- a/cpp/include/cudf/strings/detail/gather.cuh +++ b/cpp/include/cudf/strings/detail/gather.cuh @@ -112,8 +112,7 @@ std::unique_ptr gather( auto const d_out_chars = out_chars_column->mutable_view().template data(); // fill in chars - cudf::detail::device_span const d_out_offsets_span(d_out_offsets, - output_count + 1); + cudf::device_span const d_out_offsets_span(d_out_offsets, output_count + 1); auto const d_in_chars = (strings_count > 0) ? strings.chars().data() : nullptr; auto gather_chars_fn = [d_out_offsets_span, begin, d_in_offsets, d_in_chars] __device__(size_type out_char_idx) { diff --git a/cpp/include/cudf/strings/detail/strings_column_factories.cuh b/cpp/include/cudf/strings/detail/strings_column_factories.cuh index bec7a29ca18..8e843c555c5 100644 --- a/cpp/include/cudf/strings/detail/strings_column_factories.cuh +++ b/cpp/include/cudf/strings/detail/strings_column_factories.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -23,6 +23,7 @@ #include #include +#include #include #include @@ -33,7 +34,7 @@ namespace cudf { namespace strings { namespace detail { -// Create a strings-type column from vector of pointer/size pairs +// Create a strings-type column from iterators of pointer/size pairs template std::unique_ptr make_strings_column(IndexPairIterator begin, IndexPairIterator end, @@ -56,33 +57,32 @@ std::unique_ptr make_strings_column(IndexPairIterator begin, "total size of strings is too large for cudf column"); // build offsets column from the strings sizes - auto offsets_transformer = [begin] __device__(size_type idx) { - string_index_pair const item = begin[idx]; + auto offsets_transformer = [] __device__(string_index_pair item) { return (item.first != nullptr ? static_cast(item.second) : 0); }; - auto offsets_transformer_itr = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), offsets_transformer); - auto offsets_column = strings::detail::make_offsets_child_column( + auto offsets_transformer_itr = thrust::make_transform_iterator(begin, offsets_transformer); + auto offsets_column = strings::detail::make_offsets_child_column( offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr); - auto d_offsets = offsets_column->view().template data(); // create null mask auto validator = [] __device__(string_index_pair const item) { return item.first != nullptr; }; auto new_nulls = cudf::detail::valid_if(begin, end, validator, stream, mr); auto null_count = new_nulls.second; - rmm::device_buffer null_mask{0, stream, mr}; - if (null_count > 0) null_mask = std::move(new_nulls.first); + auto null_mask = + (null_count > 0) ? std::move(new_nulls.first) : rmm::device_buffer{0, stream, mr}; // build chars column auto chars_column = strings::detail::create_chars_child_column(strings_count, null_count, bytes, stream, mr); auto d_chars = chars_column->mutable_view().template data(); - auto copy_chars = [begin, d_offsets, d_chars] __device__(size_type idx) { - string_index_pair const item = begin[idx]; - if (item.first != nullptr) memcpy(d_chars + d_offsets[idx], item.first, item.second); + auto copy_chars = [d_chars] __device__(auto item) { + string_index_pair str = thrust::get<0>(item); + size_type offset = thrust::get<1>(item); + if (str.first != nullptr) memcpy(d_chars + offset, str.first, str.second); }; thrust::for_each_n(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), + thrust::make_zip_iterator( + thrust::make_tuple(begin, offsets_column->view().template begin())), strings_count, copy_chars); @@ -95,6 +95,50 @@ std::unique_ptr make_strings_column(IndexPairIterator begin, mr); } +// Create a strings-type column from iterators to chars, offsets, and bitmask. +template +std::unique_ptr make_strings_column(CharIterator chars_begin, + CharIterator chars_end, + OffsetIterator offsets_begin, + OffsetIterator offsets_end, + size_type null_count, + rmm::device_buffer&& null_mask, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + size_type strings_count = thrust::distance(offsets_begin, offsets_end) - 1; + size_type bytes = std::distance(chars_begin, chars_end) * sizeof(char); + if (strings_count == 0) return strings::detail::make_empty_strings_column(stream, mr); + + CUDF_EXPECTS(null_count < strings_count, "null strings column not yet supported"); + CUDF_EXPECTS(bytes >= 0, "invalid offsets data"); + + // build offsets column -- this is the number of strings + 1 + auto offsets_column = make_numeric_column( + data_type{type_id::INT32}, strings_count + 1, mask_state::UNALLOCATED, stream, mr); + auto offsets_view = offsets_column->mutable_view(); + thrust::transform(rmm::exec_policy(stream), + offsets_begin, + offsets_end, + offsets_view.data(), + [] __device__(auto offset) { return static_cast(offset); }); + + // build chars column + auto chars_column = + strings::detail::create_chars_child_column(strings_count, null_count, bytes, stream, mr); + auto chars_view = chars_column->mutable_view(); + thrust::copy(rmm::exec_policy(stream), chars_begin, chars_end, chars_view.data()); + + return make_strings_column(strings_count, + std::move(offsets_column), + std::move(chars_column), + null_count, + std::move(null_mask), + stream, + mr); +} + } // namespace detail } // namespace strings } // namespace cudf diff --git a/cpp/include/cudf/strings/detail/utilities.hpp b/cpp/include/cudf/strings/detail/utilities.hpp index 8844d2fb4b2..a5db4d55001 100644 --- a/cpp/include/cudf/strings/detail/utilities.hpp +++ b/cpp/include/cudf/strings/detail/utilities.hpp @@ -73,7 +73,7 @@ rmm::device_uvector create_string_vector_from_column( * @return Child offsets column */ std::unique_ptr child_offsets_from_string_vector( - cudf::detail::device_span strings, + cudf::device_span strings, rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); @@ -87,7 +87,7 @@ std::unique_ptr child_offsets_from_string_vector( * @return Child chars column */ std::unique_ptr child_chars_from_string_vector( - cudf::detail::device_span strings, + cudf::device_span strings, column_view const& offsets, rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); diff --git a/cpp/include/cudf/utilities/span.hpp b/cpp/include/cudf/utilities/span.hpp index 750eff56d4c..1f872a44fec 100644 --- a/cpp/include/cudf/utilities/span.hpp +++ b/cpp/include/cudf/utilities/span.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -28,10 +28,11 @@ #include namespace cudf { -namespace detail { constexpr std::size_t dynamic_extent = std::numeric_limits::max(); +namespace detail { + /** * @brief C++20 std::span with reduced feature set. */ @@ -100,6 +101,8 @@ class span_base { size_type _size; }; +} // namespace detail + // ===== host_span ================================================================================= template @@ -116,8 +119,8 @@ struct is_host_span_supported_container< // thrust::host_vector> : std::true_type { }; -template -struct host_span : public span_base> { +template +struct host_span : public cudf::detail::span_base> { using base = cudf::detail::span_base>; using base::base; @@ -155,8 +158,8 @@ struct is_device_span_supported_container< // rmm::device_uvector> : std::true_type { }; -template -struct device_span : public span_base> { +template +struct device_span : public cudf::detail::span_base> { using base = cudf::detail::span_base>; using base::base; @@ -173,5 +176,4 @@ struct device_span : public span_base> { } }; -} // namespace detail } // namespace cudf diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp index 510cab1ffe7..7667254ffbf 100644 --- a/cpp/include/cudf_test/column_wrapper.hpp +++ b/cpp/include/cudf_test/column_wrapper.hpp @@ -21,6 +21,7 @@ #include #include #include +#include #include #include #include @@ -701,7 +702,9 @@ class strings_column_wrapper : public detail::column_wrapper { std::vector offsets; auto all_valid = thrust::make_constant_iterator(true); std::tie(chars, offsets) = detail::make_chars_and_offsets(begin, end, all_valid); - wrapped = cudf::make_strings_column(chars, offsets); + auto d_chars = cudf::detail::make_device_uvector_sync(chars); + auto d_offsets = cudf::detail::make_device_uvector_sync(offsets); + wrapped = cudf::make_strings_column(d_chars, d_offsets); } /** @@ -740,8 +743,11 @@ class strings_column_wrapper : public detail::column_wrapper { std::vector chars; std::vector offsets; std::tie(chars, offsets) = detail::make_chars_and_offsets(begin, end, v); - wrapped = - cudf::make_strings_column(chars, offsets, detail::make_null_mask_vector(v, v + num_strings)); + auto null_mask = detail::make_null_mask_vector(v, v + num_strings); + auto d_chars = cudf::detail::make_device_uvector_sync(chars); + auto d_offsets = cudf::detail::make_device_uvector_sync(offsets); + auto d_bitmask = cudf::detail::make_device_uvector_sync(null_mask); + wrapped = cudf::make_strings_column(d_chars, d_offsets, d_bitmask); } /** diff --git a/cpp/include/cudf_test/iterator_utilities.hpp b/cpp/include/cudf_test/iterator_utilities.hpp index 40c275a13d3..297bcbf175c 100644 --- a/cpp/include/cudf_test/iterator_utilities.hpp +++ b/cpp/include/cudf_test/iterator_utilities.hpp @@ -66,7 +66,7 @@ static auto iterator_with_null_at(Iter index_start, Iter index_end) * and yields `true` (to mark valid rows) for all other indices. E.g. * * @code - * using host_span = cudf::detail::host_span; + * using host_span = cudf::host_span; * auto iter = iterator_with_null_at(host_span{std::vector{8,9}}); * iter[6] == true; // i.e. Valid row at index 6. * iter[7] == true; // i.e. Valid row at index 7. @@ -77,7 +77,7 @@ static auto iterator_with_null_at(Iter index_start, Iter index_end) * @param indices The indices for which the validity iterator must return `false` (i.e. null) * @return auto Validity iterator */ -static auto iterator_with_null_at(cudf::detail::host_span const& indices) +static auto iterator_with_null_at(cudf::host_span const& indices) { return iterator_with_null_at(indices.begin(), indices.end()); } diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu index 27f8306cbd7..7d43524f608 100644 --- a/cpp/src/binaryop/compiled/binary_ops.cu +++ b/cpp/src/binaryop/compiled/binary_ops.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -21,9 +21,10 @@ #include #include #include +#include #include -#include +#include #include namespace cudf { @@ -373,7 +374,7 @@ struct null_considering_binop { "Output column type should match input column type"); // Shallow copy of the resultant strings - rmm::device_vector out_col_strings(col_size); + rmm::device_uvector out_col_strings(col_size, stream); // Invalid output column strings - null rows cudf::string_view const invalid_str{nullptr, 0}; @@ -397,10 +398,10 @@ struct null_considering_binop { // Populate output column populate_out_col( - lhs_dev_view, rhs_dev_view, col_size, stream, minmax_func, out_col_strings.data().get()); + lhs_dev_view, rhs_dev_view, col_size, stream, minmax_func, out_col_strings.data()); // Create an output column with the resultant strings - out = make_strings_column(out_col_strings, invalid_str, stream, mr); + out = cudf::make_strings_column(out_col_strings, invalid_str, stream, mr); break; } diff --git a/cpp/src/io/avro/avro_gpu.cu b/cpp/src/io/avro/avro_gpu.cu index 4eefee66531..321f5ee8963 100644 --- a/cpp/src/io/avro/avro_gpu.cu +++ b/cpp/src/io/avro/avro_gpu.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -19,7 +19,7 @@ #include -using cudf::detail::device_span; +using cudf::device_span; namespace cudf { namespace io { diff --git a/cpp/src/io/avro/avro_gpu.h b/cpp/src/io/avro/avro_gpu.h index 5aac6f99a80..95b6e13d3f6 100644 --- a/cpp/src/io/avro/avro_gpu.h +++ b/cpp/src/io/avro/avro_gpu.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -59,7 +59,7 @@ struct schemadesc_s { */ void DecodeAvroColumnData(block_desc_s *blocks, schemadesc_s *schema, - cudf::detail::device_span global_dictionary, + cudf::device_span global_dictionary, const uint8_t *avro_data, uint32_t num_blocks, uint32_t schema_len, diff --git a/cpp/src/io/avro/reader_impl.cu b/cpp/src/io/avro/reader_impl.cu index b0806a9cf92..42035687750 100644 --- a/cpp/src/io/avro/reader_impl.cu +++ b/cpp/src/io/avro/reader_impl.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -32,7 +32,7 @@ #include #include -using cudf::detail::device_span; +using cudf::device_span; namespace cudf { namespace io { diff --git a/cpp/src/io/avro/reader_impl.hpp b/cpp/src/io/avro/reader_impl.hpp index 880c428b60d..22fa1aaa760 100644 --- a/cpp/src/io/avro/reader_impl.hpp +++ b/cpp/src/io/avro/reader_impl.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -97,7 +97,7 @@ class reader::impl { */ void decode_data(const rmm::device_buffer &block_data, const std::vector> &dict, - cudf::detail::device_span global_dictionary, + cudf::device_span global_dictionary, size_t num_rows, std::vector> columns, std::vector &out_buffers, diff --git a/cpp/src/io/comp/io_uncomp.h b/cpp/src/io/comp/io_uncomp.h index 55f8d20dda5..8daf73ecd0c 100644 --- a/cpp/src/io/comp/io_uncomp.h +++ b/cpp/src/io/comp/io_uncomp.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018, NVIDIA CORPORATION. + * Copyright (c) 2018-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. @@ -22,7 +22,7 @@ #include -using cudf::detail::host_span; +using cudf::host_span; namespace cudf { namespace io { diff --git a/cpp/src/io/comp/uncomp.cpp b/cpp/src/io/comp/uncomp.cpp index 68441ac4db9..d5166b76892 100644 --- a/cpp/src/io/comp/uncomp.cpp +++ b/cpp/src/io/comp/uncomp.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2020, NVIDIA CORPORATION. + * Copyright (c) 2018-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. @@ -26,7 +26,7 @@ #include // uncompress -using cudf::detail::host_span; +using cudf::host_span; namespace cudf { namespace io { diff --git a/cpp/src/io/csv/csv_gpu.cu b/cpp/src/io/csv/csv_gpu.cu index 041d1de3404..67c6a49ed28 100644 --- a/cpp/src/io/csv/csv_gpu.cu +++ b/cpp/src/io/csv/csv_gpu.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -42,7 +42,7 @@ using namespace ::cudf::io; -using cudf::detail::device_span; +using cudf::device_span; namespace cudf { namespace io { diff --git a/cpp/src/io/csv/csv_gpu.h b/cpp/src/io/csv/csv_gpu.h index d0e0698f8e7..0c36a1575d7 100644 --- a/cpp/src/io/csv/csv_gpu.h +++ b/cpp/src/io/csv/csv_gpu.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -23,7 +23,7 @@ #include -using cudf::detail::device_span; +using cudf::device_span; namespace cudf { namespace io { diff --git a/cpp/src/io/csv/reader_impl.cu b/cpp/src/io/csv/reader_impl.cu index 1e27ee39455..332e8aff7fc 100644 --- a/cpp/src/io/csv/reader_impl.cu +++ b/cpp/src/io/csv/reader_impl.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -42,8 +42,8 @@ using std::string; using std::vector; -using cudf::detail::device_span; -using cudf::detail::host_span; +using cudf::device_span; +using cudf::host_span; namespace cudf { namespace io { @@ -351,7 +351,7 @@ table_with_metadata reader::impl::read(rmm::cuda_stream_view stream) // during the conversion stage const std::string quotechar(1, opts.quotechar); const std::string dblquotechar(2, opts.quotechar); - std::unique_ptr col = make_strings_column(out_buffers[i]._strings, stream); + std::unique_ptr col = cudf::make_strings_column(out_buffers[i]._strings, stream); out_columns.emplace_back( cudf::strings::replace(col->view(), dblquotechar, quotechar, -1, mr_)); } else { diff --git a/cpp/src/io/csv/reader_impl.hpp b/cpp/src/io/csv/reader_impl.hpp index 67246165be0..2764eb0980c 100644 --- a/cpp/src/io/csv/reader_impl.hpp +++ b/cpp/src/io/csv/reader_impl.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -35,7 +35,7 @@ #include #include -using cudf::detail::host_span; +using cudf::host_span; namespace cudf { namespace io { diff --git a/cpp/src/io/json/json_gpu.cu b/cpp/src/io/json/json_gpu.cu index 7448d49e117..5efb64fd4d5 100644 --- a/cpp/src/io/json/json_gpu.cu +++ b/cpp/src/io/json/json_gpu.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -38,7 +38,7 @@ #include #include -using cudf::detail::device_span; +using cudf::device_span; namespace cudf { namespace io { diff --git a/cpp/src/io/json/json_gpu.h b/cpp/src/io/json/json_gpu.h index cbab408d2f1..fb8d7b2c7ab 100644 --- a/cpp/src/io/json/json_gpu.h +++ b/cpp/src/io/json/json_gpu.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -29,7 +29,7 @@ #include -using cudf::detail::device_span; +using cudf::device_span; namespace cudf { namespace io { diff --git a/cpp/src/io/json/reader_impl.cu b/cpp/src/io/json/reader_impl.cu index 4ae7e063b4b..5a82c9891b8 100644 --- a/cpp/src/io/json/reader_impl.cu +++ b/cpp/src/io/json/reader_impl.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -27,6 +27,7 @@ #include #include +#include #include #include #include @@ -41,7 +42,7 @@ #include -using cudf::detail::host_span; +using cudf::host_span; namespace cudf { namespace io { @@ -600,9 +601,22 @@ table_with_metadata reader::impl::convert_data_to_table(rmm::cuda_stream_view st stream.synchronize(); // postprocess columns - auto target = make_strings_column( - std::vector{'\\', '"', '\\', '\\', '\\', 't', '\\', 'r', '\\', 'b'}, {0, 2, 4, 6, 8, 10}); - auto repl = make_strings_column({'"', '\\', '\t', '\r', '\b'}, {0, 1, 2, 3, 4, 5}); + auto target_chars = std::vector{'\\', '"', '\\', '\\', '\\', 't', '\\', 'r', '\\', 'b'}; + auto target_offsets = std::vector{0, 2, 4, 6, 8, 10}; + + auto repl_chars = std::vector{'"', '\\', '\t', '\r', '\b'}; + auto repl_offsets = std::vector{0, 1, 2, 3, 4, 5}; + + auto target = make_strings_column(cudf::detail::make_device_uvector_async(target_chars, stream), + cudf::detail::make_device_uvector_async(target_offsets, stream), + {}, + 0, + stream); + auto repl = make_strings_column(cudf::detail::make_device_uvector_async(repl_chars, stream), + cudf::detail::make_device_uvector_async(repl_offsets, stream), + {}, + 0, + stream); thrust::host_vector h_valid_counts = d_valid_counts; std::vector> out_columns; @@ -619,6 +633,10 @@ table_with_metadata reader::impl::convert_data_to_table(rmm::cuda_stream_view st } } + // This is to ensure the stream-ordered make_stream_column calls above complete before + // the temporary std::vectors are destroyed on exit from this function. + stream.synchronize(); + CUDF_EXPECTS(!out_columns.empty(), "No columns created from json input"); return table_with_metadata{std::make_unique(std::move(out_columns)), metadata_}; diff --git a/cpp/src/io/orc/timezone.cuh b/cpp/src/io/orc/timezone.cuh index 6bb1e787432..3a87f28391c 100644 --- a/cpp/src/io/orc/timezone.cuh +++ b/cpp/src/io/orc/timezone.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -32,8 +32,8 @@ namespace io { struct timezone_table_view { int32_t gmt_offset = 0; - cudf::detail::device_span ttimes; - cudf::detail::device_span offsets; + cudf::device_span ttimes; + cudf::device_span offsets; }; static constexpr int64_t day_seconds = 24 * 60 * 60; @@ -85,8 +85,8 @@ CUDA_HOST_DEVICE_CALLABLE int32_t get_gmt_offset_impl(int64_t const *ttimes, * * Implemented in `get_gmt_offset_impl`. */ -inline __host__ int32_t get_gmt_offset(cudf::detail::host_span ttimes, - cudf::detail::host_span offsets, +inline __host__ int32_t get_gmt_offset(cudf::host_span ttimes, + cudf::host_span offsets, int64_t ts) { CUDF_EXPECTS(ttimes.size() == offsets.size(), @@ -99,8 +99,8 @@ inline __host__ int32_t get_gmt_offset(cudf::detail::host_span tt * * Implemented in `get_gmt_offset_impl`. */ -inline __device__ int32_t get_gmt_offset(cudf::detail::device_span ttimes, - cudf::detail::device_span offsets, +inline __device__ int32_t get_gmt_offset(cudf::device_span ttimes, + cudf::device_span offsets, int64_t ts) { return get_gmt_offset_impl(ttimes.begin(), offsets.begin(), ttimes.size(), ts); diff --git a/cpp/src/io/utilities/parsing_utils.cuh b/cpp/src/io/utilities/parsing_utils.cuh index c7f405e1cc0..584d2c9a74a 100644 --- a/cpp/src/io/utilities/parsing_utils.cuh +++ b/cpp/src/io/utilities/parsing_utils.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -22,7 +22,7 @@ #include -using cudf::detail::device_span; +using cudf::device_span; namespace cudf { namespace io { diff --git a/cpp/src/reductions/scan.cu b/cpp/src/reductions/scan.cu index 011b34031fe..f73ffb0214a 100644 --- a/cpp/src/reductions/scan.cu +++ b/cpp/src/reductions/scan.cu @@ -25,10 +25,11 @@ #include #include #include +#include #include #include -#include +#include #include namespace cudf { @@ -166,22 +167,21 @@ struct scan_dispatcher { rmm::mr::device_memory_resource* mr) { const size_type size = input_view.size(); - rmm::device_vector result(size); + rmm::device_uvector result(size, stream); auto d_input = column_device_view::create(input_view, stream); if (input_view.has_nulls()) { auto input = make_null_replacement_iterator(*d_input, Op::template identity()); - thrust::inclusive_scan( - rmm::exec_policy(stream), input, input + size, result.data().get(), Op{}); + thrust::inclusive_scan(rmm::exec_policy(stream), input, input + size, result.data(), Op{}); } else { auto input = d_input->begin(); - thrust::inclusive_scan( - rmm::exec_policy(stream), input, input + size, result.data().get(), Op{}); + thrust::inclusive_scan(rmm::exec_policy(stream), input, input + size, result.data(), Op{}); } CHECK_CUDA(stream.value()); - auto output_column = make_strings_column(result, Op::template identity(), stream, mr); + auto output_column = + cudf::make_strings_column(result, Op::template identity(), stream, mr); if (null_handling == null_policy::EXCLUDE) { output_column->set_null_mask(detail::copy_bitmask(input_view, stream, mr), input_view.null_count()); diff --git a/cpp/src/strings/convert/convert_urls.cu b/cpp/src/strings/convert/convert_urls.cu index 2b93995ec87..cdca23a3584 100644 --- a/cpp/src/strings/convert/convert_urls.cu +++ b/cpp/src/strings/convert/convert_urls.cu @@ -33,7 +33,7 @@ #include #include -using cudf::detail::device_span; +using cudf::device_span; namespace cudf { namespace strings { diff --git a/cpp/src/strings/extract.cu b/cpp/src/strings/extract.cu index e2bb43dc783..f33c0c01fb6 100644 --- a/cpp/src/strings/extract.cu +++ b/cpp/src/strings/extract.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -92,26 +92,25 @@ std::unique_ptr
extract( auto regex_insts = d_prog.insts_counts(); for (int32_t column_index = 0; column_index < groups; ++column_index) { - rmm::device_vector indices(strings_count); - string_index_pair* d_indices = indices.data().get(); + rmm::device_uvector indices(strings_count, stream); if ((regex_insts > MAX_STACK_INSTS) || (regex_insts <= RX_SMALL_INSTS)) thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(strings_count), - d_indices, + indices.begin(), extract_fn{d_prog, d_strings, column_index}); else if (regex_insts <= RX_MEDIUM_INSTS) thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(strings_count), - d_indices, + indices.begin(), extract_fn{d_prog, d_strings, column_index}); else thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(strings_count), - d_indices, + indices.begin(), extract_fn{d_prog, d_strings, column_index}); results.emplace_back(make_strings_column(indices, stream, mr)); diff --git a/cpp/src/strings/findall.cu b/cpp/src/strings/findall.cu index 2c26875b5d6..bcd9c808271 100644 --- a/cpp/src/strings/findall.cu +++ b/cpp/src/strings/findall.cu @@ -166,28 +166,27 @@ std::unique_ptr
findall_re( for (int32_t column_index = 0; column_index < columns; ++column_index) { rmm::device_uvector indices(strings_count, stream); - auto d_indices = indices.data(); if ((regex_insts > MAX_STACK_INSTS) || (regex_insts <= RX_SMALL_INSTS)) thrust::transform( rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(strings_count), - d_indices, + indices.begin(), findall_fn{*d_strings, *d_prog, column_index, d_find_counts}); else if (regex_insts <= RX_MEDIUM_INSTS) thrust::transform( rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(strings_count), - d_indices, + indices.begin(), findall_fn{*d_strings, *d_prog, column_index, d_find_counts}); else thrust::transform( rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(strings_count), - d_indices, + indices.begin(), findall_fn{*d_strings, *d_prog, column_index, d_find_counts}); // results.emplace_back(make_strings_column(indices.begin(), indices.end(), stream, mr)); diff --git a/cpp/src/strings/replace/replace.cu b/cpp/src/strings/replace/replace.cu index 75d8b626409..0955c217526 100644 --- a/cpp/src/strings/replace/replace.cu +++ b/cpp/src/strings/replace/replace.cu @@ -38,8 +38,6 @@ #include #include -using cudf::detail::device_span; - namespace cudf { namespace strings { namespace detail { diff --git a/cpp/src/strings/split/partition.cu b/cpp/src/strings/split/partition.cu index b6c636557d6..aa096f60333 100644 --- a/cpp/src/strings/split/partition.cu +++ b/cpp/src/strings/split/partition.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -60,14 +60,14 @@ struct partition_fn { partition_fn(column_device_view const& d_strings, string_view const& d_delimiter, - rmm::device_vector& indices_left, - rmm::device_vector& indices_delim, - rmm::device_vector& indices_right) + rmm::device_uvector& indices_left, + rmm::device_uvector& indices_delim, + rmm::device_uvector& indices_right) : d_strings(d_strings), d_delimiter(d_delimiter), - d_indices_left(indices_left.data().get()), - d_indices_delim(indices_delim.data().get()), - d_indices_right(indices_right.data().get()) + d_indices_left(indices_left.data()), + d_indices_delim(indices_delim.data()), + d_indices_right(indices_right.data()) { } @@ -145,9 +145,9 @@ struct partition_fn { struct rpartition_fn : public partition_fn { rpartition_fn(column_device_view const& d_strings, string_view const& d_delimiter, - rmm::device_vector& indices_left, - rmm::device_vector& indices_delim, - rmm::device_vector& indices_right) + rmm::device_uvector& indices_left, + rmm::device_uvector& indices_delim, + rmm::device_uvector& indices_right) : partition_fn(d_strings, d_delimiter, indices_left, indices_delim, indices_right) { } @@ -187,8 +187,9 @@ std::unique_ptr
partition( if (strings_count == 0) return std::make_unique
(std::vector>()); auto strings_column = column_device_view::create(strings.parent(), stream); string_view d_delimiter(delimiter.data(), delimiter.size()); - rmm::device_vector left_indices(strings_count), delim_indices(strings_count), - right_indices(strings_count); + auto left_indices = rmm::device_uvector(strings_count, stream); + auto delim_indices = rmm::device_uvector(strings_count, stream); + auto right_indices = rmm::device_uvector(strings_count, stream); partition_fn partitioner( *strings_column, d_delimiter, left_indices, delim_indices, right_indices); @@ -214,8 +215,9 @@ std::unique_ptr
rpartition( if (strings_count == 0) return std::make_unique
(std::vector>()); auto strings_column = column_device_view::create(strings.parent(), stream); string_view d_delimiter(delimiter.data(), delimiter.size()); - rmm::device_vector left_indices(strings_count), delim_indices(strings_count), - right_indices(strings_count); + auto left_indices = rmm::device_uvector(strings_count, stream); + auto delim_indices = rmm::device_uvector(strings_count, stream); + auto right_indices = rmm::device_uvector(strings_count, stream); rpartition_fn partitioner( *strings_column, d_delimiter, left_indices, delim_indices, right_indices); thrust::for_each_n(rmm::exec_policy(stream), diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu index 774c6fa0da6..4d6c9389173 100644 --- a/cpp/src/strings/strings_column_factories.cu +++ b/cpp/src/strings/strings_column_factories.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -17,93 +17,18 @@ #include #include #include -#include -#include +#include #include +#include #include #include #include #include -#include -#include - namespace cudf { -// Create a strings-type column from vector of pointer/size pairs -std::unique_ptr make_strings_column( - const rmm::device_vector>& strings, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - CUDF_FUNC_RANGE(); - size_type strings_count = strings.size(); - if (strings_count == 0) return strings::detail::make_empty_strings_column(stream, mr); - - auto d_strings = strings.data().get(); - - // check total size is not too large for cudf column - auto size_checker = [d_strings] __device__(size_t idx) { - auto item = d_strings[idx]; - return (item.first != nullptr) ? item.second : 0; - }; - size_t bytes = thrust::transform_reduce(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(strings_count), - size_checker, - 0, - thrust::plus()); - CUDF_EXPECTS(bytes < static_cast(std::numeric_limits::max()), - "total size of strings is too large for cudf column"); - - // build offsets column from the strings sizes - auto offsets_transformer = [d_strings] __device__(size_type idx) { - thrust::pair item = d_strings[idx]; - return (item.first != nullptr ? static_cast(item.second) : 0); - }; - auto offsets_transformer_itr = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), offsets_transformer); - auto offsets_column = strings::detail::make_offsets_child_column( - offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr); - auto offsets_view = offsets_column->view(); - auto d_offsets = offsets_view.data(); - - // create null mask - auto new_nulls = detail::valid_if( - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(strings_count), - [d_strings] __device__(size_type idx) { return d_strings[idx].first != nullptr; }, - stream, - mr); - auto null_count = new_nulls.second; - rmm::device_buffer null_mask{0, stream, mr}; - if (null_count > 0) null_mask = std::move(new_nulls.first); - - // build chars column - auto chars_column = - strings::detail::create_chars_child_column(strings_count, null_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_offsets, d_chars] __device__(size_type idx) { - // place individual strings - auto item = d_strings[idx]; - if (item.first != nullptr) - memcpy(d_chars + d_offsets[idx], item.first, item.second); - }); - - return make_strings_column(strings_count, - std::move(offsets_column), - std::move(chars_column), - null_count, - std::move(null_mask), - stream, - mr); -} - +namespace { struct string_view_to_pair { string_view null_placeholder; string_view_to_pair(string_view n) : null_placeholder(n) {} @@ -115,89 +40,74 @@ struct string_view_to_pair { } }; -// Create a strings-type column from vector of string_view -std::unique_ptr make_strings_column(const rmm::device_vector& string_views, - const string_view null_placeholder, +} // namespace + +// Create a strings-type column from vector of pointer/size pairs +std::unique_ptr make_strings_column( + device_span const> strings, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + + return cudf::strings::detail::make_strings_column(strings.begin(), strings.end(), stream, mr); +} + +std::unique_ptr make_strings_column( + device_span chars, + device_span offsets, + size_type null_count, + rmm::device_buffer&& null_mask, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + CUDF_FUNC_RANGE(); + + return cudf::strings::detail::make_strings_column(chars.begin(), + chars.end(), + offsets.begin(), + offsets.end(), + null_count, + std::move(null_mask), + stream, + mr); +} + +std::unique_ptr make_strings_column(device_span string_views, + string_view null_placeholder, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { + CUDF_FUNC_RANGE(); + auto it_pair = thrust::make_transform_iterator(string_views.begin(), string_view_to_pair{null_placeholder}); - const rmm::device_vector> dev_strings( - it_pair, it_pair + string_views.size()); - return make_strings_column(dev_strings, stream, mr); + return cudf::strings::detail::make_strings_column( + it_pair, it_pair + string_views.size(), stream, mr); } // Create a strings-type column from device vector of chars and vector of offsets. -std::unique_ptr make_strings_column(const rmm::device_vector& strings, - const rmm::device_vector& offsets, - const rmm::device_vector& valid_mask, +std::unique_ptr make_strings_column(cudf::device_span strings, + cudf::device_span offsets, + cudf::device_span valid_mask, size_type null_count, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - size_type num_strings = offsets.size() - 1; - if (num_strings == 0) return strings::detail::make_empty_strings_column(stream, mr); - - CUDF_EXPECTS(null_count < num_strings, "null strings column not yet supported"); - if (null_count > 0) { - CUDF_EXPECTS(!valid_mask.empty(), "Cannot have null elements without a null mask."); - } - size_type bytes = offsets.back(); - CUDF_EXPECTS(bytes >= 0, "invalid offsets vector"); - - // build offsets column -- this is the number of strings + 1 - auto offsets_column = make_numeric_column( - data_type{type_id::INT32}, num_strings + 1, mask_state::UNALLOCATED, stream, mr); - auto offsets_view = offsets_column->mutable_view(); - thrust::transform(rmm::exec_policy(stream), - offsets.begin(), - offsets.end(), - offsets_view.data(), - [] __device__(auto offset) { return static_cast(offset); }); // build null bitmask rmm::device_buffer null_mask{ - valid_mask.data().get(), - valid_mask.size() * - sizeof(bitmask_type)}; // Or this works too: sizeof(typename - // std::remove_reference_t::value_type) - // Following give the incorrect value of 8 instead of 4 because of smart references: - // sizeof(valid_mask[0]), sizeof(decltype(valid_mask.front())) - - // build chars column - auto chars_column = - strings::detail::create_chars_child_column(num_strings, null_count, bytes, stream, mr); - auto chars_view = chars_column->mutable_view(); - CUDA_TRY(cudaMemcpyAsync(chars_view.data(), - strings.data().get(), - bytes, - cudaMemcpyDeviceToDevice, - stream.value())); - - return make_strings_column(num_strings, - std::move(offsets_column), - std::move(chars_column), - null_count, - std::move(null_mask), - stream, - mr); -} - -// Create strings column from host vectors -std::unique_ptr make_strings_column(const std::vector& strings, - const std::vector& offsets, - const std::vector& null_mask, - size_type null_count, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - rmm::device_vector d_strings{strings}; - rmm::device_vector d_offsets{offsets}; - rmm::device_vector d_null_mask{null_mask}; - - return make_strings_column(d_strings, d_offsets, d_null_mask, null_count, stream, mr); + valid_mask.data(), valid_mask.size() * sizeof(bitmask_type), stream, mr}; + + return cudf::strings::detail::make_strings_column(strings.begin(), + strings.end(), + offsets.begin(), + offsets.end(), + null_count, + std::move(null_mask), + stream, + mr); } // @@ -209,6 +119,8 @@ std::unique_ptr make_strings_column(size_type num_strings, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { + CUDF_FUNC_RANGE(); + if (null_count > 0) CUDF_EXPECTS(null_mask.size() > 0, "Column with nulls must be nullable."); CUDF_EXPECTS(num_strings == offsets_column->size() - 1, "Invalid offsets column size for strings column."); diff --git a/cpp/src/strings/utilities.cu b/cpp/src/strings/utilities.cu index 5b9a1374224..2af313627ad 100644 --- a/cpp/src/strings/utilities.cu +++ b/cpp/src/strings/utilities.cu @@ -80,7 +80,7 @@ rmm::device_uvector create_string_vector_from_column(cudf::strings_ * @copydoc child_offsets_from_string_vector */ std::unique_ptr child_offsets_from_string_vector( - cudf::detail::device_span strings, + cudf::device_span strings, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { @@ -90,11 +90,10 @@ std::unique_ptr child_offsets_from_string_vector( /** * @copydoc child_chars_from_string_vector */ -std::unique_ptr child_chars_from_string_vector( - cudf::detail::device_span strings, - column_view const& offsets, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::unique_ptr child_chars_from_string_vector(cudf::device_span strings, + column_view const& offsets, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { auto const d_strings = strings.data(); auto const bytes = cudf::detail::get_value(offsets, strings.size(), stream); diff --git a/cpp/src/text/tokenize.cu b/cpp/src/text/tokenize.cu index 438f3bff4b7..0ba51f7639f 100644 --- a/cpp/src/text/tokenize.cu +++ b/cpp/src/text/tokenize.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -70,18 +70,18 @@ std::unique_ptr tokenize_fn(cudf::size_type strings_count, token_count_fn(strings_count, tokenizer, stream, rmm::mr::get_current_device_resource()); auto d_token_counts = token_counts->view(); // create token-index offsets from the counts - rmm::device_vector token_offsets(strings_count + 1); + rmm::device_uvector token_offsets(strings_count + 1, stream); thrust::inclusive_scan(rmm::exec_policy(stream), d_token_counts.template begin(), d_token_counts.template end(), token_offsets.begin() + 1); - CUDA_TRY(cudaMemsetAsync(token_offsets.data().get(), 0, sizeof(int32_t), stream.value())); - auto const total_tokens = token_offsets.back(); + CUDA_TRY(cudaMemsetAsync(token_offsets.data(), 0, sizeof(int32_t), stream.value())); + auto const total_tokens = token_offsets.back_element(stream); // build a list of pointers to each token - rmm::device_vector tokens(total_tokens); + rmm::device_uvector tokens(total_tokens, stream); // now go get the tokens - tokenizer.d_offsets = token_offsets.data().get(); - tokenizer.d_tokens = tokens.data().get(); + tokenizer.d_offsets = token_offsets.data(); + tokenizer.d_tokens = tokens.data(); thrust::for_each_n(rmm::exec_policy(stream), thrust::make_counting_iterator(0), strings_count, diff --git a/cpp/tests/strings/factories_test.cu b/cpp/tests/strings/factories_test.cu index 810be3d6d60..f904c404251 100644 --- a/cpp/tests/strings/factories_test.cu +++ b/cpp/tests/strings/factories_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019, NVIDIA CORPORATION. + * 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. @@ -146,10 +146,6 @@ TEST_F(StringsFactoriesTest, CreateColumnFromOffsets) EXPECT_EQ(memcmp(h_buffer.data(), h_chars_data.data(), h_buffer.size()), 0); EXPECT_EQ( memcmp(h_offsets.data(), h_offsets_data.data(), h_offsets.size() * sizeof(cudf::size_type)), 0); - - // check host version of the factory too - auto column2 = cudf::make_strings_column(h_buffer, h_offsets, h_nulls, null_count); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(column->view(), column2->view()); } TEST_F(StringsFactoriesTest, CreateScalar) diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index 7a775056b3f..cea66eced11 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -190,7 +190,8 @@ std::string stringify_column_differences(thrust::device_vector const& diffe fixed_width_column_wrapper(h_differences.begin(), h_differences.end()); auto diff_table = cudf::gather(source_table, diff_column); // Need to pull back the differences - auto const h_left_strings = to_strings(diff_table->get_column(0)); + auto const h_left_strings = to_strings(diff_table->get_column(0)); + auto const h_right_strings = to_strings(diff_table->get_column(1)); for (size_t i = 0; i < h_differences.size(); ++i) buffer << depth_str << "lhs[" << h_differences[i] << "] = " << h_left_strings[i] << ", rhs[" diff --git a/cpp/tests/utilities_tests/span_tests.cu b/cpp/tests/utilities_tests/span_tests.cu index d49a345fc61..547894e9f6c 100644 --- a/cpp/tests/utilities_tests/span_tests.cu +++ b/cpp/tests/utilities_tests/span_tests.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -27,8 +27,8 @@ #include #include -using cudf::detail::device_span; -using cudf::detail::host_span; +using cudf::device_span; +using cudf::host_span; template void expect_equivolent(host_span a, host_span b)