diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 3066fdf32e77..f558d91b7ddc 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -581,6 +581,10 @@ add_library( src/ast/expressions.cpp src/ast/jit/expressions.cpp src/ast/operators.cpp + src/binary/binary_column_factories.cpp + src/binary/binary_column_view.cpp + src/binary/copying.cu + src/binary/from_views.cu src/binaryop/binaryop.cpp src/binaryop/compiled/ATan2.cu src/binaryop/compiled/Add.cu diff --git a/cpp/include/cudf/binary/binary_column_factories.hpp b/cpp/include/cudf/binary/binary_column_factories.hpp new file mode 100644 index 000000000000..00838dc7f977 --- /dev/null +++ b/cpp/include/cudf/binary/binary_column_factories.hpp @@ -0,0 +1,69 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include +#include +#include +#include +#include + +#include + +#include + +/** + * @file + * @brief Factory functions for `BINARY` columns. + */ + +namespace CUDF_EXPORT cudf { + +/** + * @brief Constructs a `BINARY` column by copying device-resident binary views. + * + * A view whose data pointer equals `null_placeholder.data()` becomes null. + * + * @param binary_views Device span of binary values + * @param null_placeholder View whose data pointer marks null values + * @param stream CUDA stream used for device operations + * @param mr Device memory resource used for output allocations + * @return Newly constructed `BINARY` column + */ +std::unique_ptr make_binary_column( + device_span binary_views, + binary_view null_placeholder, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + +/** + * @brief Constructs a `BINARY` column from offsets and a contiguous byte buffer. + * + * `offsets_column` must contain either `INT32` or `INT64` values and must not + * contain nulls. Nonnegative, monotonically nondecreasing offsets whose final + * value does not exceed `bytes_buffer.size()` are a caller precondition. + * + * Empty input is canonicalized to an empty `BINARY` column with no children. + * + * @param num_rows Number of binary values represented by the column + * @param offsets_column Offsets with `num_rows + 1` elements + * @param bytes_buffer Contiguous payload bytes + * @param null_count Number of null rows + * @param null_mask Row validity mask + * @return Newly constructed `BINARY` column + */ +std::unique_ptr make_binary_column(size_type num_rows, + std::unique_ptr offsets_column, + rmm::device_buffer&& bytes_buffer, + size_type null_count, + rmm::device_buffer&& null_mask); + +/** + * @brief Creates an empty `BINARY` column. + */ +std::unique_ptr make_empty_binary_column(); + +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/binary/binary_column_view.hpp b/cpp/include/cudf/binary/binary_column_view.hpp new file mode 100644 index 000000000000..d52ca9f5c3ad --- /dev/null +++ b/cpp/include/cudf/binary/binary_column_view.hpp @@ -0,0 +1,57 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include + +#include + +#include + +/** + * @file + * @brief Class definition for `cudf::binary_column_view`. + */ + +namespace CUDF_EXPORT cudf { + +/** + * @brief A wrapper providing access to a `BINARY` column's offsets and payload. + */ +class binary_column_view : private column_view { + public: + static constexpr size_type offsets_column_index{0}; + + explicit binary_column_view(column_view binary_column); + binary_column_view() = default; + binary_column_view(binary_column_view const&) = default; + binary_column_view(binary_column_view&&) = default; + ~binary_column_view() override = default; + binary_column_view& operator=(binary_column_view const&) = default; + binary_column_view& operator=(binary_column_view&&) = default; + + using column_view::has_nulls; + using column_view::is_empty; + using column_view::null_count; + using column_view::null_mask; + using column_view::offset; + using column_view::size; + + [[nodiscard]] column_view parent() const; + [[nodiscard]] column_view offsets() const; + + /** + * @brief Returns the total number of bytes in the underlying payload buffer. + * + * This reports the unsliced parent's payload size. + */ + [[nodiscard]] int64_t bytes_size(rmm::cuda_stream_view stream) const; + + [[nodiscard]] uint8_t const* bytes_begin() const noexcept; + [[nodiscard]] uint8_t const* bytes_end(rmm::cuda_stream_view stream) const; +}; + +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/binary/binary_view.hpp b/cpp/include/cudf/binary/binary_view.hpp new file mode 100644 index 000000000000..f432bc55d745 --- /dev/null +++ b/cpp/include/cudf/binary/binary_view.hpp @@ -0,0 +1,123 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include + +#include + +#include + +/** + * @file + * @brief Class definition for `cudf::binary_view`. + */ + +namespace CUDF_EXPORT cudf { + +/** + * @brief A non-owning, immutable view of a variable-length sequence of bytes. + * + * The caller must maintain the device memory for the lifetime of this object. + * Binary values have no character encoding; comparisons use unsigned-byte + * lexicographic ordering. + */ +class binary_view { + public: + using value_type = uint8_t; + using const_iterator = value_type const*; + + CUDF_HOST_DEVICE constexpr binary_view() = default; + + /** + * @brief Constructs a view over `size` bytes beginning at `data`. + */ + CUDF_HOST_DEVICE constexpr binary_view(value_type const* data, size_type size) + : _data{data}, _size{size} + { + } + + /** + * @brief Returns the number of bytes in this value. + */ + CUDF_HOST_DEVICE [[nodiscard]] constexpr size_type size_bytes() const noexcept { return _size; } + + /** + * @brief Returns a pointer to the first byte. + */ + CUDF_HOST_DEVICE [[nodiscard]] constexpr value_type const* data() const noexcept + { + return _data; + } + + /** + * @brief Returns whether this value contains no bytes. + */ + CUDF_HOST_DEVICE [[nodiscard]] constexpr bool empty() const noexcept { return _size == 0; } + + CUDF_HOST_DEVICE [[nodiscard]] constexpr const_iterator begin() const noexcept { return _data; } + CUDF_HOST_DEVICE [[nodiscard]] constexpr const_iterator end() const noexcept + { + return _data + _size; + } + + /** + * @brief Returns the byte at `index`. + */ + CUDF_HOST_DEVICE [[nodiscard]] constexpr value_type operator[](size_type index) const noexcept + { + return _data[index]; + } + + /** + * @brief Compares values using unsigned-byte lexicographic ordering. + */ + CUDF_HOST_DEVICE [[nodiscard]] constexpr int compare(binary_view rhs) const noexcept + { + auto const count = cuda::std::min(_size, rhs._size); + for (size_type i = 0; i < count; ++i) { + if (_data[i] < rhs._data[i]) { return -1; } + if (_data[i] > rhs._data[i]) { return 1; } + } + return (_size > rhs._size) - (_size < rhs._size); + } + + CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator==(binary_view rhs) const noexcept + { + return compare(rhs) == 0; + } + + CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator!=(binary_view rhs) const noexcept + { + return not(*this == rhs); + } + + CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator<(binary_view rhs) const noexcept + { + return compare(rhs) < 0; + } + + CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator>(binary_view rhs) const noexcept + { + return compare(rhs) > 0; + } + + CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator<=(binary_view rhs) const noexcept + { + return compare(rhs) <= 0; + } + + CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator>=(binary_view rhs) const noexcept + { + return compare(rhs) >= 0; + } + + private: + value_type const* _data{}; + size_type _size{}; +}; + +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/binary/detail/binary_column_factories.cuh b/cpp/include/cudf/binary/detail/binary_column_factories.cuh new file mode 100644 index 000000000000..55e886a36b72 --- /dev/null +++ b/cpp/include/cudf/binary/detail/binary_column_factories.cuh @@ -0,0 +1,66 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include +#include +#include + +#include +#include +#include +#include + +namespace cudf::binary::detail { + +using binary_index_pair = cuda::std::pair; + +/** + * @brief Creates a BINARY column from device-accessible pointer/size pairs. + * + * A null pointer denotes a null row. The bytes referenced by non-null pointers + * are copied into the output payload. + */ +template +std::unique_ptr make_binary_column(IndexPairIterator begin, + IndexPairIterator end, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + auto const row_count = static_cast(cuda::std::distance(begin, end)); + if (row_count == 0) { return make_empty_binary_column(); } + + auto sizes = thrust::make_transform_iterator( + begin, + cuda::proclaim_return_type([] __device__(binary_index_pair item) { + return item.first == nullptr ? size_type{0} : item.second; + })); + auto [offsets, bytes] = cudf::strings::detail::make_offsets_child_column( + sizes, sizes + row_count, stream, mr); + + auto const validator = [] __device__(binary_index_pair item) { return item.first != nullptr; }; + auto [null_mask, null_count] = cudf::detail::valid_if(begin, end, validator, stream, mr); + if (null_count == 0) { null_mask = rmm::device_buffer{0, stream, mr}; } + + auto char_pairs = thrust::make_transform_iterator( + begin, + cuda::proclaim_return_type>( + [] __device__(binary_index_pair item) { + return cuda::std::pair{ + reinterpret_cast(item.first), item.second}; + })); + auto payload = cudf::strings::detail::make_chars_buffer( + offsets->view(), bytes, char_pairs, row_count, stream, mr); + + return cudf::make_binary_column(row_count, + std::move(offsets), + payload.release(), + null_count, + std::move(null_mask)); +} + +} // namespace cudf::binary::detail diff --git a/cpp/include/cudf/binary/detail/copying.hpp b/cpp/include/cudf/binary/detail/copying.hpp new file mode 100644 index 000000000000..f3b5257865e3 --- /dev/null +++ b/cpp/include/cudf/binary/detail/copying.hpp @@ -0,0 +1,31 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include +#include +#include + +#include + +#include + +namespace cudf::binary::detail { + +/** + * @brief Copies rows `[start, end)` from a BINARY column into a new owning column. + * + * The output offsets are normalized to begin at zero and retain the input + * offsets width. + */ +CUDF_EXPORT std::unique_ptr copy_slice( + binary_column_view const& input, + size_type start, + size_type end, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + +} // namespace cudf::binary::detail diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh index a7add0c42c26..d4babd7e0698 100644 --- a/cpp/include/cudf/column/column_device_view.cuh +++ b/cpp/include/cudf/column/column_device_view.cuh @@ -4,6 +4,7 @@ */ #pragma once +#include #include #include #include @@ -132,7 +133,9 @@ class alignas(16) column_device_view : public column_device_view_core { * @param element_index Position of the desired string element * @return string_view instance representing this element at this index */ - template )> + template or + cuda::std::is_same_v)> [[nodiscard]] __device__ T element(size_type element_index) const noexcept { return base::element(element_index); diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh index dcbf3b95292d..ed0dfc49d7f0 100644 --- a/cpp/include/cudf/column/column_device_view_base.cuh +++ b/cpp/include/cudf/column/column_device_view_base.cuh @@ -4,6 +4,7 @@ */ #pragma once +#include #include #include #include @@ -464,15 +465,21 @@ class alignas(16) column_device_view_core : public detail::column_device_view_ba * @param element_index Position of the desired string element * @return string_view instance representing this element at this index */ - template )> + template or + cuda::std::is_same_v)> [[nodiscard]] __device__ T element(size_type element_index) const noexcept { - size_type index = element_index + offset(); // account for this view's _offset - char const* d_strings = static_cast(_data); - auto const offsets = child(offsets_column_index); - auto const itr = cudf::detail::input_offsetalator(offsets.head(), offsets.type()); - auto const offset = itr[index]; - return string_view{d_strings + offset, static_cast(itr[index + 1] - offset)}; + size_type index = element_index + offset(); // account for this view's _offset + auto const offsets = child(offsets_column_index); + auto const itr = cudf::detail::input_offsetalator(offsets.head(), offsets.type()); + auto const offset = itr[index]; + auto const size = static_cast(itr[index + 1] - offset); + if constexpr (cuda::std::is_same_v) { + return string_view{static_cast(_data) + offset, size}; + } else { + return binary_view{static_cast(_data) + offset, size}; + } } public: @@ -697,15 +704,21 @@ class alignas(16) mutable_column_device_view_core : public detail::column_device * @param element_index Position of the desired string element * @return string_view instance representing this element at this index */ - template )> + template or + cuda::std::is_same_v)> [[nodiscard]] __device__ T element(size_type element_index) const noexcept { - size_type index = element_index + offset(); // account for this view's _offset - char const* d_strings = static_cast(_data); - auto const offsets = child(offsets_column_index); - auto const itr = cudf::detail::input_offsetalator(offsets.head(), offsets.type()); - auto const offset = itr[index]; - return string_view{d_strings + offset, static_cast(itr[index + 1] - offset)}; + size_type index = element_index + offset(); // account for this view's _offset + auto const offsets = child(offsets_column_index); + auto const itr = cudf::detail::input_offsetalator(offsets.head(), offsets.type()); + auto const offset = itr[index]; + auto const size = static_cast(itr[index + 1] - offset); + if constexpr (cuda::std::is_same_v) { + return string_view{static_cast(_data) + offset, size}; + } else { + return binary_view{static_cast(_data) + offset, size}; + } } /** diff --git a/cpp/include/cudf/detail/gather.cuh b/cpp/include/cudf/detail/gather.cuh index e5bb1f9ff575..48593b0e7594 100644 --- a/cpp/include/cudf/detail/gather.cuh +++ b/cpp/include/cudf/detail/gather.cuh @@ -4,6 +4,7 @@ */ #pragma once +#include #include #include #include @@ -28,6 +29,7 @@ #include #include +#include #include #include @@ -264,6 +266,48 @@ struct column_gatherer_impl { } }; +/** + * @brief Column gather specialization for BINARY columns. + */ +template <> +struct column_gatherer_impl { + template + std::unique_ptr operator()(column_view const& source_column, + MapIterator gather_map_begin, + MapIterator gather_map_end, + bool nullify_out_of_bounds, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) + { + auto const row_count = + static_cast(cudf::distance(gather_map_begin, gather_map_end)); + if (row_count == 0) { return make_empty_binary_column(); } + + using map_type = typename std::iterator_traits::value_type; + auto source = column_device_view::create(source_column, stream); + auto values = rmm::device_uvector(row_count, stream, mr); + thrust::transform( + rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + gather_map_begin, + gather_map_end, + values.begin(), + [source = *source, nullify_out_of_bounds] __device__(auto index) { + if (nullify_out_of_bounds && + !bounds_checker{0, source.size()}(static_cast(index))) { + return binary_view{}; + } + if (source.is_null(index)) { return binary_view{}; } + auto const value = source.element(index); + auto const data = value.data() == nullptr && value.empty() + ? reinterpret_cast(1) + : value.data(); + return binary_view{data, value.size_bytes()}; + }); + return make_binary_column( + device_span{values}, binary_view{}, stream, mr); + } +}; + /** * @brief Column gather specialization for list_view column type. * diff --git a/cpp/include/cudf/detail/scatter.cuh b/cpp/include/cudf/detail/scatter.cuh index a836d7eb33eb..ed560b4134d6 100644 --- a/cpp/include/cudf/detail/scatter.cuh +++ b/cpp/include/cudf/detail/scatter.cuh @@ -172,6 +172,57 @@ struct column_scatterer_impl { } }; +template <> +struct column_scatterer_impl { + template + std::unique_ptr operator()(column_view const& source, + MapIterator scatter_map_begin, + MapIterator scatter_map_end, + column_view const& target, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) const + { + auto values = rmm::device_uvector(target.size(), stream, mr); + auto target_device = column_device_view::create(target, stream); + auto const target_has_nulls = target.has_nulls(); + thrust::transform( + rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator{0}, + cuda::counting_iterator{target.size()}, + values.begin(), + [column = *target_device, target_has_nulls] __device__(size_type index) { + if (target_has_nulls && column.is_null(index)) { return binary_view{}; } + auto const value = column.element(index); + auto const data = value.data() == nullptr && value.empty() + ? reinterpret_cast(1) + : value.data(); + return binary_view{data, value.size_bytes()}; + }); + + auto source_device = column_device_view::create(source, stream); + auto const source_has_nulls = source.has_nulls(); + auto source_values = thrust::make_transform_iterator( + cuda::counting_iterator{0}, + cuda::proclaim_return_type( + [column = *source_device, source_has_nulls] __device__(size_type index) { + if (source_has_nulls && column.is_null(index)) { return binary_view{}; } + auto const value = column.element(index); + auto const data = value.data() == nullptr && value.empty() + ? reinterpret_cast(1) + : value.data(); + return binary_view{data, value.size_bytes()}; + })); + thrust::scatter(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + source_values, + source_values + cudf::distance(scatter_map_begin, scatter_map_end), + scatter_map_begin, + values.begin()); + + return make_binary_column( + device_span{values}, binary_view{}, stream, mr); + } +}; + template <> struct column_scatterer_impl { template diff --git a/cpp/include/cudf/hashing/detail/murmurhash3_x64_128.cuh b/cpp/include/cudf/hashing/detail/murmurhash3_x64_128.cuh index 15d62c3d15b2..5f864bb3d686 100644 --- a/cpp/include/cudf/hashing/detail/murmurhash3_x64_128.cuh +++ b/cpp/include/cudf/hashing/detail/murmurhash3_x64_128.cuh @@ -4,6 +4,7 @@ */ #pragma once +#include #include #include #include @@ -72,6 +73,15 @@ MurmurHash3_x64_128::result_type key.size_bytes()); } +template <> +MurmurHash3_x64_128::result_type + __device__ inline MurmurHash3_x64_128:: + operator()(cudf::binary_view const& key) const +{ + return this->compute_bytes(reinterpret_cast(key.data()), + key.size_bytes()); +} + template <> MurmurHash3_x64_128::result_type __device__ inline MurmurHash3_x64_128:: diff --git a/cpp/include/cudf/hashing/detail/murmurhash3_x86_32.cuh b/cpp/include/cudf/hashing/detail/murmurhash3_x86_32.cuh index dbae072947be..145e7abc2610 100644 --- a/cpp/include/cudf/hashing/detail/murmurhash3_x86_32.cuh +++ b/cpp/include/cudf/hashing/detail/murmurhash3_x86_32.cuh @@ -5,6 +5,7 @@ #pragma once +#include #include #include #include @@ -75,6 +76,15 @@ MurmurHash3_x86_32::result_type key.size_bytes()); } +template <> +MurmurHash3_x86_32::result_type + __device__ inline MurmurHash3_x86_32:: + operator()(cudf::binary_view const& key) const +{ + return this->compute_bytes(reinterpret_cast(key.data()), + key.size_bytes()); +} + template <> MurmurHash3_x86_32::result_type __device__ inline MurmurHash3_x86_32:: diff --git a/cpp/include/cudf/hashing/detail/xxhash_32.cuh b/cpp/include/cudf/hashing/detail/xxhash_32.cuh index 99cf810cde22..444c3408697e 100644 --- a/cpp/include/cudf/hashing/detail/xxhash_32.cuh +++ b/cpp/include/cudf/hashing/detail/xxhash_32.cuh @@ -5,6 +5,7 @@ #pragma once +#include #include #include #include @@ -69,6 +70,14 @@ operator()(cudf::string_view const& key) const key.size_bytes()); } +template <> +XXHash_32::result_type __device__ inline XXHash_32:: +operator()(cudf::binary_view const& key) const +{ + return this->compute_bytes(reinterpret_cast(key.data()), + key.size_bytes()); +} + template <> XXHash_32::result_type __device__ inline XXHash_32:: operator()(numeric::decimal32 const& key) const diff --git a/cpp/include/cudf/hashing/detail/xxhash_64.cuh b/cpp/include/cudf/hashing/detail/xxhash_64.cuh index eb918589278d..d904178e0122 100644 --- a/cpp/include/cudf/hashing/detail/xxhash_64.cuh +++ b/cpp/include/cudf/hashing/detail/xxhash_64.cuh @@ -5,6 +5,7 @@ #pragma once +#include #include #include #include @@ -67,6 +68,14 @@ operator()(cudf::string_view const& key) const key.size_bytes()); } +template <> +XXHash_64::result_type __device__ inline XXHash_64:: +operator()(cudf::binary_view const& key) const +{ + return this->compute_bytes(reinterpret_cast(key.data()), + key.size_bytes()); +} + template <> XXHash_64::result_type __device__ inline XXHash_64:: operator()(numeric::decimal32 const& key) const diff --git a/cpp/include/cudf/io/types.hpp b/cpp/include/cudf/io/types.hpp index 639b7bd9d316..2654f2b60480 100644 --- a/cpp/include/cudf/io/types.hpp +++ b/cpp/include/cudf/io/types.hpp @@ -961,12 +961,20 @@ struct partition_info { }; /** - * @brief schema element for reader - * + * @brief Output representation for Parquet byte-array data. + */ +enum class byte_array_output_type : int8_t { + DEFAULT, ///< Use the file schema: annotated text becomes STRING, otherwise BINARY + BINARY, ///< Return a BINARY column + STRING, ///< Return a STRING column + LIST_UINT8, ///< Return the legacy LIST representation +}; + +/** + * @brief Schema element for reader output options. */ class reader_column_schema { - // Whether to read binary data as a string column - bool _convert_binary_to_strings{true}; + byte_array_output_type _byte_array_output{byte_array_output_type::DEFAULT}; int32_t _type_length{0}; std::vector children; @@ -1020,16 +1028,27 @@ class reader_column_schema { [[nodiscard]] reader_column_schema const& child(size_type i) const { return children[i]; } /** - * @brief Specifies whether this column should be written as binary or string data - * Only valid for the following column types: - * string, list + * @brief Compatibility API selecting STRING or legacy LIST output. * - * @param convert_to_string True = convert binary to strings False = return binary + * @param convert_to_string True selects STRING; false selects LIST * @return this for chaining */ reader_column_schema& set_convert_binary_to_strings(bool convert_to_string) { - _convert_binary_to_strings = convert_to_string; + _byte_array_output = + convert_to_string ? byte_array_output_type::STRING : byte_array_output_type::LIST_UINT8; + return *this; + } + + /** + * @brief Selects the output representation for Parquet byte-array data. + * + * @param output_type Requested output representation + * @return this for chaining + */ + reader_column_schema& set_byte_array_output(byte_array_output_type output_type) + { + _byte_array_output = output_type; return *this; } @@ -1046,13 +1065,21 @@ class reader_column_schema { } /** - * @brief Get whether to encode this column as binary or string data + * @brief Returns whether STRING output was explicitly selected. * - * @return Boolean indicating whether to encode this column as binary data + * @return true when the byte-array output policy is STRING */ [[nodiscard]] bool is_enabled_convert_binary_to_strings() const { - return _convert_binary_to_strings; + return _byte_array_output == byte_array_output_type::STRING; + } + + /** + * @brief Returns the requested output representation for byte-array data. + */ + [[nodiscard]] byte_array_output_type get_byte_array_output() const + { + return _byte_array_output; } /** diff --git a/cpp/include/cudf/types.hpp b/cpp/include/cudf/types.hpp index cc3b30c98e8b..f466e92e84bf 100644 --- a/cpp/include/cudf/types.hpp +++ b/cpp/include/cudf/types.hpp @@ -43,6 +43,7 @@ namespace CUDF_EXPORT cudf { class column; class column_view; class mutable_column_view; +class binary_view; class string_view; class list_view; class struct_view; @@ -212,6 +213,7 @@ enum class type_id : int32_t { DECIMAL64, ///< Fixed-point type with int64_t DECIMAL128, ///< Fixed-point type with __int128_t STRUCT, ///< Struct elements + BINARY, ///< Variable-length byte sequence elements // `NUM_TYPE_IDS` must be last! NUM_TYPE_IDS ///< Total number of type ids }; diff --git a/cpp/include/cudf/utilities/traits.hpp b/cpp/include/cudf/utilities/traits.hpp index 6f35d05d6f04..94988508a2d3 100644 --- a/cpp/include/cudf/utilities/traits.hpp +++ b/cpp/include/cudf/utilities/traits.hpp @@ -599,6 +599,7 @@ CUDF_HOST_DEVICE constexpr inline bool is_fixed_width() */ bool is_fixed_width(data_type type); +class binary_view; class string_view; /** @@ -616,7 +617,8 @@ class string_view; template CUDF_HOST_DEVICE constexpr inline bool is_compound() { - return cuda::std::is_same_v or + return cuda::std::is_same_v or + cuda::std::is_same_v or cuda::std::is_same_v or cuda::std::is_same_v or cuda::std::is_same_v; } diff --git a/cpp/include/cudf/utilities/type_dispatcher.hpp b/cpp/include/cudf/utilities/type_dispatcher.hpp index 074cee275a39..49a9eb18ff8d 100644 --- a/cpp/include/cudf/utilities/type_dispatcher.hpp +++ b/cpp/include/cudf/utilities/type_dispatcher.hpp @@ -176,6 +176,7 @@ CUDF_TYPE_MAPPING(numeric::decimal32, type_id::DECIMAL32) CUDF_TYPE_MAPPING(numeric::decimal64, type_id::DECIMAL64) CUDF_TYPE_MAPPING(numeric::decimal128, type_id::DECIMAL128) CUDF_TYPE_MAPPING(cudf::struct_view, type_id::STRUCT) +CUDF_TYPE_MAPPING(cudf::binary_view, type_id::BINARY) /** * @brief Specialization to map 'char' type to type_id::INT8 @@ -550,6 +551,9 @@ CUDF_HOST_DEVICE __forceinline__ constexpr decltype(auto) type_dispatcher(cudf:: case type_id::STRUCT: return f.template operator()::type>( std::forward(args)...); + case type_id::BINARY: + return f.template operator()::type>( + std::forward(args)...); default: { #ifndef __CUDA_ARCH__ CUDF_FAIL("Invalid type_id."); diff --git a/cpp/src/binary/binary_column_factories.cpp b/cpp/src/binary/binary_column_factories.cpp new file mode 100644 index 000000000000..e60d78a8ef99 --- /dev/null +++ b/cpp/src/binary/binary_column_factories.cpp @@ -0,0 +1,56 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include + +#include + +namespace cudf { + +std::unique_ptr make_empty_binary_column() +{ + return make_empty_column(data_type{type_id::BINARY}); +} + +std::unique_ptr make_binary_column(size_type num_rows, + std::unique_ptr offsets_column, + rmm::device_buffer&& bytes_buffer, + size_type null_count, + rmm::device_buffer&& null_mask) +{ + CUDF_EXPECTS(num_rows >= 0, "Number of binary rows cannot be negative"); + CUDF_EXPECTS(offsets_column != nullptr, "Offsets column must not be null"); + CUDF_EXPECTS(offsets_column->type().id() == type_id::INT32 or + offsets_column->type().id() == type_id::INT64, + "Binary offsets must have type INT32 or INT64"); + CUDF_EXPECTS(offsets_column->null_count() == 0, "Binary offsets must not contain nulls"); + CUDF_EXPECTS(null_count >= 0 and null_count <= num_rows, "Invalid binary null count"); + if (null_count > 0) { + CUDF_EXPECTS(null_mask.size() > 0, "Binary column with nulls must be nullable"); + } + + if (num_rows == 0) { + CUDF_EXPECTS(offsets_column->size() == 0 or offsets_column->size() == 1, + "Empty binary offsets must contain zero or one element"); + CUDF_EXPECTS(bytes_buffer.size() == 0, "Empty binary column cannot contain payload bytes"); + return make_empty_binary_column(); + } + + CUDF_EXPECTS(offsets_column->size() == num_rows + 1, + "Binary offsets size must equal the row count plus one"); + + std::vector> children; + children.emplace_back(std::move(offsets_column)); + return std::make_unique(data_type{type_id::BINARY}, + num_rows, + std::move(bytes_buffer), + std::move(null_mask), + null_count, + std::move(children)); +} + +} // namespace cudf diff --git a/cpp/src/binary/binary_column_view.cpp b/cpp/src/binary/binary_column_view.cpp new file mode 100644 index 000000000000..cf75208887f6 --- /dev/null +++ b/cpp/src/binary/binary_column_view.cpp @@ -0,0 +1,55 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include + +namespace cudf { + +binary_column_view::binary_column_view(column_view binary_column) : column_view(binary_column) +{ + CUDF_EXPECTS(type().id() == type_id::BINARY, "binary_column_view only supports BINARY"); + if (not is_empty()) { + CUDF_EXPECTS(num_children() == 1, "non-empty binary column must have one offsets child"); + auto const offsets_view = offsets(); + CUDF_EXPECTS(offsets_view.type().id() == type_id::INT32 or + offsets_view.type().id() == type_id::INT64, + "binary offsets must have type INT32 or INT64"); + CUDF_EXPECTS(offsets_view.null_count() == 0, "binary offsets must not contain nulls"); + CUDF_EXPECTS(offsets_view.size() >= offset() + size() + 1, + "binary offsets do not cover the column view"); + } +} + +column_view binary_column_view::parent() const { return static_cast(*this); } + +column_view binary_column_view::offsets() const +{ + CUDF_EXPECTS(num_children() > 0, "binary column has no children"); + return child(offsets_column_index); +} + +int64_t binary_column_view::bytes_size(rmm::cuda_stream_view stream) const +{ + if (size() == 0) { return 0; } + auto const offsets_view = offsets(); + auto const last = offsets_view.size() - 1; + if (offsets_view.type().id() == type_id::INT32) { + return cudf::detail::get_value(offsets_view, last, stream); + } + CUDF_EXPECTS(offsets_view.type().id() == type_id::INT64, + "binary offsets must have type INT32 or INT64"); + return cudf::detail::get_value(offsets_view, last, stream); +} + +uint8_t const* binary_column_view::bytes_begin() const noexcept { return head(); } + +uint8_t const* binary_column_view::bytes_end(rmm::cuda_stream_view stream) const +{ + return bytes_begin() + bytes_size(stream); +} + +} // namespace cudf diff --git a/cpp/src/binary/copying.cu b/cpp/src/binary/copying.cu new file mode 100644 index 000000000000..8e68ca858de6 --- /dev/null +++ b/cpp/src/binary/copying.cu @@ -0,0 +1,84 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +#include +#include + +namespace cudf::binary::detail { +namespace { + +int64_t get_offset(column_view const& offsets, size_type index, rmm::cuda_stream_view stream) +{ + if (offsets.type().id() == type_id::INT32) { + return cudf::detail::get_value(offsets, index, stream); + } + CUDF_EXPECTS(offsets.type().id() == type_id::INT64, + "Binary offsets must have type INT32 or INT64"); + return cudf::detail::get_value(offsets, index, stream); +} + +} // namespace + +std::unique_ptr copy_slice(binary_column_view const& input, + size_type start, + size_type end, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_EXPECTS(start >= 0 and start <= end and end <= input.size(), + "Invalid BINARY slice range"); + if (start == end) { return make_empty_binary_column(); } + + auto const row_count = end - start; + auto const offsets_offset = input.offset() + start; + auto offsets_column = std::make_unique( + cudf::detail::slice( + input.offsets(), {offsets_offset, offsets_offset + row_count + 1}, stream) + .front(), + stream, + mr); + + auto const first_offset = get_offset(offsets_column->view(), 0, stream); + if (first_offset != 0) { + auto output_offsets = + cudf::detail::offsetalator_factory::make_output_iterator(offsets_column->mutable_view()); + auto input_offsets = + cudf::detail::offsetalator_factory::make_input_iterator(input.offsets(), offsets_offset); + thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + input_offsets, + input_offsets + offsets_column->size(), + output_offsets, + cuda::proclaim_return_type( + [first_offset] __device__(auto offset) { return offset - first_offset; })); + } + + auto const payload_size = + static_cast(get_offset(offsets_column->view(), row_count, stream)); + auto payload = rmm::device_buffer{input.bytes_begin() + first_offset, payload_size, stream, mr}; + + auto null_mask = cudf::detail::copy_bitmask( + input.null_mask(), offsets_offset, offsets_offset + row_count, stream, mr); + auto const null_count = cudf::detail::null_count( + static_cast(null_mask.data()), 0, row_count, stream); + + return make_binary_column(row_count, + std::move(offsets_column), + std::move(payload), + null_count, + std::move(null_mask)); +} + +} // namespace cudf::binary::detail diff --git a/cpp/src/binary/from_views.cu b/cpp/src/binary/from_views.cu new file mode 100644 index 000000000000..0540c053af58 --- /dev/null +++ b/cpp/src/binary/from_views.cu @@ -0,0 +1,38 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include + +#include +#include + +namespace cudf { +namespace { + +struct binary_view_to_pair { + binary_view null_placeholder; + + __device__ binary::detail::binary_index_pair operator()(binary_view value) const + { + return value.data() == null_placeholder.data() + ? binary::detail::binary_index_pair{nullptr, 0} + : binary::detail::binary_index_pair{value.data(), value.size_bytes()}; + } +}; + +} // namespace + +std::unique_ptr make_binary_column(device_span binary_views, + binary_view null_placeholder, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto begin = + thrust::make_transform_iterator(binary_views.begin(), binary_view_to_pair{null_placeholder}); + return binary::detail::make_binary_column(begin, begin + binary_views.size(), stream, mr); +} + +} // namespace cudf diff --git a/cpp/src/column/column.cu b/cpp/src/column/column.cu index 085fe0d9562c..86955b5f4738 100644 --- a/cpp/src/column/column.cu +++ b/cpp/src/column/column.cu @@ -3,6 +3,8 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include +#include #include #include #include @@ -168,6 +170,14 @@ struct create_column_from_view { return cudf::strings::detail::copy_slice(sview, 0, view.size(), stream, mr); } + template + std::unique_ptr operator()() + requires(std::is_same_v) + { + return cudf::binary::detail::copy_slice( + cudf::binary_column_view{view}, 0, view.size(), stream, mr); + } + template std::unique_ptr operator()() requires(std::is_same_v) diff --git a/cpp/src/column/column_view.cpp b/cpp/src/column/column_view.cpp index b9973b4a2eab..a50f12392c5a 100644 --- a/cpp/src/column/column_view.cpp +++ b/cpp/src/column/column_view.cpp @@ -4,6 +4,7 @@ */ #include +#include #include #include #include @@ -44,6 +45,11 @@ void prefetch_col_data(ColumnView& col, void const* data_ptr, std::string_view k data_ptr, scv.chars_size(cudf::get_default_stream()) * sizeof(char), cudf::get_default_stream()); + } else if (col.type().id() == type_id::BINARY) { + binary_column_view const bcv{col}; + if (data_ptr == nullptr) { return; } + cudf::prefetch::detail::prefetch_noexcept( + data_ptr, bcv.bytes_size(cudf::get_default_stream()), cudf::get_default_stream()); } else { CUDF_LOG_DEBUG("Unsupported type: %d", static_cast(col.type().id())); } @@ -119,7 +125,7 @@ column_view_base::column_view_base(data_type type, CUDF_EXPECTS(nullptr == data, "EMPTY column should have no data."); CUDF_EXPECTS(nullptr == null_mask, "EMPTY column should have no null mask."); } else if (is_compound(type)) { - if (type.id() != type_id::STRING) { + if (type.id() != type_id::STRING and type.id() != type_id::BINARY) { CUDF_EXPECTS(nullptr == data, "Compound (parent) columns cannot have data"); } } else if (size > 0) { diff --git a/cpp/src/copying/concatenate.cu b/cpp/src/copying/concatenate.cu index 4ebd84072da3..6b8263e251d2 100644 --- a/cpp/src/copying/concatenate.cu +++ b/cpp/src/copying/concatenate.cu @@ -3,6 +3,7 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include #include #include #include @@ -27,6 +28,7 @@ #include #include +#include #include #include @@ -360,6 +362,41 @@ std::unique_ptr concatenate_dispatch::operator()() return cudf::strings::detail::concatenate(views, stream, mr); } +template <> +std::unique_ptr concatenate_dispatch::operator()() +{ + auto const row_count = std::accumulate( + views.begin(), views.end(), size_type{0}, [](size_type total, column_view const& view) { + return total + view.size(); + }); + if (row_count == 0) { return make_empty_binary_column(); } + + auto values = rmm::device_uvector(row_count, stream, mr); + size_type position = 0; + for (auto const& view : views) { + if (view.is_empty()) { continue; } + auto device_view = column_device_view::create(view, stream); + auto const has_nulls = view.has_nulls(); + auto input = cuda::counting_iterator{0}; + thrust::transform( + rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + input, + input + view.size(), + values.begin() + position, + [column = *device_view, has_nulls] __device__(size_type index) { + if (has_nulls && column.is_null(index)) { return binary_view{}; } + auto const value = column.element(index); + auto const data = value.data() == nullptr && value.empty() + ? reinterpret_cast(1) + : value.data(); + return binary_view{data, value.size_bytes()}; + }); + position += view.size(); + } + return make_binary_column( + device_span{values}, binary_view{}, stream, mr); +} + template <> std::unique_ptr concatenate_dispatch::operator()() { diff --git a/cpp/src/copying/get_element.cu b/cpp/src/copying/get_element.cu index 3ff512d15465..6952511e41d9 100644 --- a/cpp/src/copying/get_element.cu +++ b/cpp/src/copying/get_element.cu @@ -134,6 +134,15 @@ struct get_element_functor { } } + template >* p = nullptr> + std::unique_ptr operator()(column_view const&, + size_type, + cuda::stream_ref, + rmm::device_async_resource_ref) + { + CUDF_FAIL("BINARY scalar support is not yet implemented"); + } + template ()>* p = nullptr> std::unique_ptr operator()(column_view const& input, size_type index, diff --git a/cpp/src/copying/scatter.cu b/cpp/src/copying/scatter.cu index c8b10f15740c..9f81c46ffb9f 100644 --- a/cpp/src/copying/scatter.cu +++ b/cpp/src/copying/scatter.cu @@ -150,6 +150,19 @@ struct column_scalar_scatterer_impl { } }; +template +struct column_scalar_scatterer_impl { + std::unique_ptr operator()(std::reference_wrapper const&, + MapIterator, + size_type, + column_view const&, + cuda::stream_ref, + rmm::device_async_resource_ref) const + { + CUDF_FAIL("Scattering a BINARY scalar is not yet supported"); + } +}; + template struct column_scalar_scatterer_impl { std::unique_ptr operator()(std::reference_wrapper const& source, diff --git a/cpp/src/dictionary/search.cu b/cpp/src/dictionary/search.cu index d0e868572907..dc64800dc6fa 100644 --- a/cpp/src/dictionary/search.cu +++ b/cpp/src/dictionary/search.cu @@ -38,7 +38,8 @@ struct find_index_fn { cuda::stream_ref stream, rmm::device_async_resource_ref mr) const requires(not std::is_same_v and - not std::is_same_v and not std::is_same_v) + not std::is_same_v and not std::is_same_v and + not std::is_same_v) { auto const num_keys = input.keys_size(); if (!key.is_valid(stream) || num_keys == 0) { @@ -72,11 +73,10 @@ struct find_index_fn { scalar const&, cuda::stream_ref, rmm::device_async_resource_ref) const - requires(std::is_same_v or std::is_same_v or - std::is_same_v) + requires(std::is_same_v or std::is_same_v or + std::is_same_v or std::is_same_v) { - CUDF_FAIL( - "dictionary, list_view, and struct_view columns cannot be the keys column of a dictionary"); + CUDF_FAIL("BINARY, dictionary, LIST, and STRUCT columns cannot currently be dictionary keys"); } }; diff --git a/cpp/src/groupby/sort/group_scan_util.cuh b/cpp/src/groupby/sort/group_scan_util.cuh index 922a7f939432..a31a0ad23025 100644 --- a/cpp/src/groupby/sort/group_scan_util.cuh +++ b/cpp/src/groupby/sort/group_scan_util.cuh @@ -68,7 +68,7 @@ static constexpr bool is_group_scan_supported() else if (K == aggregation::PRODUCT) return cudf::is_numeric(); else if (K == aggregation::MIN or K == aggregation::MAX) - return not cudf::is_dictionary() and + return not cudf::is_dictionary() and not std::is_same_v and (is_relationally_comparable() or std::is_same_v); else return false; diff --git a/cpp/src/io/json/parser_features.cpp b/cpp/src/io/json/parser_features.cpp index 68869579b5da..661823fdc790 100644 --- a/cpp/src/io/json/parser_features.cpp +++ b/cpp/src/io/json/parser_features.cpp @@ -5,6 +5,7 @@ #include "nested_json.hpp" +#include #include #include #include @@ -149,6 +150,16 @@ struct allnull_column_functor { return make_strings_column( size, std::move(offsets), rmm::device_buffer{}, size, std::move(null_mask)); } + + template )> + std::unique_ptr operator()(schema_element const&, size_type size) const + { + auto offsets = make_zeroed_offsets(size); + auto null_mask = cudf::detail::create_null_mask(size, mask_state::ALL_NULL, stream, mr); + return make_binary_column( + size, std::move(offsets), rmm::device_buffer{}, size, std::move(null_mask)); + } + template )> std::unique_ptr operator()(schema_element const& schema, size_type size) const { diff --git a/cpp/src/io/parquet/arrow_schema_writer.cpp b/cpp/src/io/parquet/arrow_schema_writer.cpp index 9e883324d5c0..7324ebd7352e 100644 --- a/cpp/src/io/parquet/arrow_schema_writer.cpp +++ b/cpp/src/io/parquet/arrow_schema_writer.cpp @@ -162,6 +162,14 @@ struct dispatch_to_flatbuf { field_offset = flatbuf::CreateUtf8View(fbb).Union(); } + template + void operator()() + requires(std::is_same_v) + { + field_type_id = flatbuf::Type_BinaryView; + field_offset = flatbuf::CreateBinaryView(fbb).Union(); + } + template void operator()() requires(std::is_same_v) diff --git a/cpp/src/io/parquet/chunk_dict.cu b/cpp/src/io/parquet/chunk_dict.cu index dddd82c5b303..c3153b6dec6b 100644 --- a/cpp/src/io/parquet/chunk_dict.cu +++ b/cpp/src/io/parquet/chunk_dict.cu @@ -5,6 +5,7 @@ #include "parquet_gpu.cuh" +#include #include #include #include @@ -153,14 +154,16 @@ struct map_insert_fn { if (col_type == type_id::STRING) { // Strings are stored as int32_t length + string bytes return sizeof(int32_t) + data_col.element(val_idx).size_bytes(); + } else if (col_type == type_id::BINARY) { + return sizeof(int32_t) + data_col.element(val_idx).size_bytes(); } else if (col_type == type_id::LIST) { // Binary is stored as int32_t length + bytes return sizeof(int32_t) + get_element(data_col, val_idx).size_bytes(); } CUDF_UNREACHABLE( - "Byte array only supports string and list column types for dictionary " - "encoding!"); + "Byte array only supports string, binary, and list column types for " + "dictionary encoding!"); } case Type::FIXED_LEN_BYTE_ARRAY: if (data_col.type().id() == type_id::DECIMAL128) { return sizeof(__int128_t); } diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index aab25c0e648b..4692f7dffd12 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -1022,12 +1022,8 @@ table_with_metadata hybrid_scan_reader_impl::read_chunk_internal( : std::nullopt; auto const& schema = _extended_metadata->get_schema(_output_column_schemas[i]); auto const logical_type = schema.logical_type.value_or(LogicalType{}); - // FIXED_LEN_BYTE_ARRAY never read as string. - // TODO: if we ever decide that the default reader behavior is to treat unannotated BINARY - // as binary and not strings, this test needs to change. if (schema.type == Type::FIXED_LEN_BYTE_ARRAY and logical_type.type != LogicalType::DECIMAL) { - metadata = std::make_optional(); - metadata->set_convert_binary_to_strings(false); + if (not metadata.has_value()) { metadata.emplace(); } metadata->set_type_length(schema.type_length); } // Only construct `out_metadata` if `_output_metadata` has not been cached. diff --git a/cpp/src/io/parquet/page_enc.cu b/cpp/src/io/parquet/page_enc.cu index 4c5c5523313c..f7b5c77d6d1f 100644 --- a/cpp/src/io/parquet/page_enc.cu +++ b/cpp/src/io/parquet/page_enc.cu @@ -9,6 +9,7 @@ #include "page_string_utils.cuh" #include "parquet_gpu.cuh" +#include #include #include #include @@ -199,6 +200,10 @@ void __device__ calculate_frag_size(frag_init_state_s* const s, int t) auto str = s->col.leaf_column->element(val_idx); len += str.size_bytes(); } break; + case type_id::BINARY: { + auto value = s->col.leaf_column->element(val_idx); + len += value.size_bytes(); + } break; case type_id::LIST: { auto list_element = get_element(*s->col.leaf_column, val_idx); @@ -1722,6 +1727,8 @@ CUDF_KERNEL void __launch_bounds__(block_size, 8) if (physical_type == Type::BYTE_ARRAY) { if (type_id == type_id::STRING) { len += s->col.leaf_column->element(val_idx).size_bytes(); + } else if (type_id == type_id::BINARY) { + len += s->col.leaf_column->element(val_idx).size_bytes(); } else if (s->col.output_as_byte_array && type_id == type_id::LIST) { len += get_element(*s->col.leaf_column, val_idx).size_bytes(); @@ -1815,6 +1822,9 @@ CUDF_KERNEL void __launch_bounds__(block_size, 8) case type_id::STRING: return reinterpret_cast( leaf_column->element(val_idx).data()); + case type_id::BINARY: + return reinterpret_cast( + leaf_column->element(val_idx).data()); case type_id::LIST: return reinterpret_cast( get_element(*(leaf_column), val_idx).data()); @@ -1841,8 +1851,15 @@ CUDF_KERNEL void __launch_bounds__(block_size, 8) dst + pos); } } else { - auto const elem = - get_element(*(s->col.leaf_column), val_idx); + auto const elem = [&] { + if (type_id == type_id::BINARY) { + auto const value = s->col.leaf_column->element(val_idx); + return statistics::byte_array_view{ + reinterpret_cast(value.data()), + static_cast(value.size_bytes())}; + } + return get_element(*(s->col.leaf_column), val_idx); + }(); if (len != 0 and elem.data() != nullptr) { if (is_split_stream) { auto const v_char_ptr = reinterpret_cast(elem.data()); @@ -2166,6 +2183,8 @@ CUDF_KERNEL void __launch_bounds__(block_size, 8) if (type_id == type_id::STRING) { first_string = reinterpret_cast( s->col.leaf_column->element(idx_in_col).data()); + } else if (type_id == type_id::BINARY) { + first_string = s->col.leaf_column->element(idx_in_col).data(); } else if (s->col.output_as_byte_array && type_id == type_id::LIST) { first_string = reinterpret_cast( get_element(*s->col.leaf_column, idx_in_col).data()); @@ -2195,6 +2214,8 @@ CUDF_KERNEL void __launch_bounds__(block_size, 8) if (is_valid) { if (type_id == type_id::STRING) { v = s->col.leaf_column->element(val_idx).size_bytes(); + } else if (type_id == type_id::BINARY) { + v = s->col.leaf_column->element(val_idx).size_bytes(); } else if (s->col.output_as_byte_array && type_id == type_id::LIST) { auto const arr_size = get_element(*s->col.leaf_column, val_idx).size_bytes(); @@ -2357,6 +2378,9 @@ CUDF_KERNEL void __launch_bounds__(block_size, 8) if (type_id == type_id::STRING) { auto const str = s->col.leaf_column->element(idx); return {reinterpret_cast(str.data()), str.size_bytes()}; + } else if (type_id == type_id::BINARY) { + auto const value = s->col.leaf_column->element(idx); + return {value.data(), value.size_bytes()}; } else if (s->col.output_as_byte_array && type_id == type_id::LIST) { auto const str = get_element(*s->col.leaf_column, idx); return {reinterpret_cast(str.data()), diff --git a/cpp/src/io/parquet/reader_impl.cpp b/cpp/src/io/parquet/reader_impl.cpp index 75394523b594..5ad3ba4d455c 100644 --- a/cpp/src/io/parquet/reader_impl.cpp +++ b/cpp/src/io/parquet/reader_impl.cpp @@ -750,12 +750,8 @@ table_with_metadata reader_impl::read_chunk_internal(read_mode mode) : std::nullopt; auto const& schema = _metadata->get_schema(_output_column_schemas[i]); auto const logical_type = schema.logical_type.value_or(LogicalType{}); - // FIXED_LEN_BYTE_ARRAY never read as string. - // TODO: if we ever decide that the default reader behavior is to treat unannotated BINARY as - // binary and not strings, this test needs to change. if (schema.type == Type::FIXED_LEN_BYTE_ARRAY and logical_type.type != LogicalType::DECIMAL) { - metadata = std::make_optional(); - metadata->set_convert_binary_to_strings(false); + if (not metadata.has_value()) { metadata.emplace(); } metadata->set_type_length(schema.type_length); } // Only construct `out_metadata` if `_output_metadata` has not been cached. diff --git a/cpp/src/io/parquet/reader_impl_helpers.cpp b/cpp/src/io/parquet/reader_impl_helpers.cpp index b91f54cc3001..a3fe5e736ac4 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.cpp +++ b/cpp/src/io/parquet/reader_impl_helpers.cpp @@ -132,6 +132,29 @@ cuda::std::optional converted_to_logical_type(SchemaElement const& return cuda::std::nullopt; } +bool has_string_annotation(SchemaElement const& schema) +{ + if (schema.logical_type.has_value()) { + switch (schema.logical_type->type) { + case LogicalType::STRING: + case LogicalType::ENUM: + case LogicalType::JSON: + case LogicalType::BSON: return true; + default: break; + } + } + if (schema.converted_type.has_value()) { + switch (*schema.converted_type) { + case ConvertedType::UTF8: + case ConvertedType::ENUM: + case ConvertedType::JSON: + case ConvertedType::BSON: return true; + default: break; + } + } + return false; +} + /** * @brief Lookup schema children by name or field ID across multiple sources. */ @@ -1956,12 +1979,12 @@ aggregate_reader_metadata::select_columns( // pop off the extra nesting element. if (one_level_list) { nesting.pop_back(); } - // Flag the `metadata` / `value` BYTE_ARRAY children of a VARIANT group so that - // `make_column` materializes them as `list` instead of strings. - if (schema_elem.type == Type::BYTE_ARRAY && parent_schema.logical_type.has_value() && - parent_schema.logical_type->type == LogicalType::VARIANT && - (schema_elem.name == "metadata" || schema_elem.name == "value")) { - output_col.string_as_binary = true; + if (output_col.type.id() == type_id::STRING && + (schema_elem.type == Type::BYTE_ARRAY || + schema_elem.type == Type::FIXED_LEN_BYTE_ARRAY)) { + output_col.byte_array_output = has_string_annotation(schema_elem) + ? byte_array_output_type::STRING + : byte_array_output_type::BINARY; } path_is_valid = true; // If we're able to reach leaf then path is valid diff --git a/cpp/src/io/parquet/writer_impl.cu b/cpp/src/io/parquet/writer_impl.cu index 30d8c2171e56..5b468fbd9403 100644 --- a/cpp/src/io/parquet/writer_impl.cu +++ b/cpp/src/io/parquet/writer_impl.cu @@ -21,6 +21,7 @@ #include "writer_impl.hpp" #include "writer_impl_helpers.hpp" +#include #include #include #include @@ -271,11 +272,13 @@ size_t column_size(column_view const& column, rmm::cuda_stream_view stream) if (is_fixed_width(column.type())) { return size_of(column.type()) * column.size(); - } else if (column.type().id() == type_id::STRING) { - auto const scol = strings_column_view(column); + } else if (column.type().id() == type_id::STRING or column.type().id() == type_id::BINARY) { + auto const offsets = column.type().id() == type_id::STRING + ? strings_column_view(column).offsets() + : binary_column_view(column).offsets(); return cudf::strings::detail::get_offset_value( - scol.offsets(), column.size() + column.offset(), stream) - - cudf::strings::detail::get_offset_value(scol.offsets(), column.offset(), stream); + offsets, column.size() + column.offset(), stream) - + cudf::strings::detail::get_offset_value(offsets, column.offset(), stream); } else if (column.type().id() == type_id::STRUCT) { auto const scol = structs_column_view(column); size_t ret = 0; @@ -443,6 +446,16 @@ struct leaf_schema_fn { } } + template + void operator()() + requires(std::is_same_v) + { + col_schema.type = + col_meta.get_type_length() > 0 ? Type::FIXED_LEN_BYTE_ARRAY : Type::BYTE_ARRAY; + col_schema.type_length = col_meta.get_type_length(); + col_schema.stats_dtype = statistics_dtype::dtype_byte_array; + } + template void operator()() requires(std::is_same_v) @@ -898,13 +911,13 @@ std::vector construct_parquet_schema_tree( } else { // if leaf, add current - if (col->type().id() == type_id::STRING) { - if (col_meta.is_enabled_output_as_binary()) { + if (col->type().id() == type_id::STRING or col->type().id() == type_id::BINARY) { + if (col->type().id() == type_id::STRING and col_meta.is_enabled_output_as_binary()) { CUDF_EXPECTS(col_meta.num_children() == 2 or col_meta.num_children() == 0, "Binary column's corresponding metadata should have zero or two children"); } else { CUDF_EXPECTS(col_meta.num_children() == 1 or col_meta.num_children() == 0, - "String column's corresponding metadata should have zero or one children"); + "Variable-width leaf column metadata should have zero or one children"); } } else { CUDF_EXPECTS(col_meta.num_children() == 0, diff --git a/cpp/src/io/parquet/writer_impl_helpers.cpp b/cpp/src/io/parquet/writer_impl_helpers.cpp index a32f1be2243f..0b3fbfa58e4b 100644 --- a/cpp/src/io/parquet/writer_impl_helpers.cpp +++ b/cpp/src/io/parquet/writer_impl_helpers.cpp @@ -10,6 +10,7 @@ #include "writer_impl_helpers.hpp" +#include #include #include #include @@ -43,11 +44,13 @@ void fill_table_meta(table_input_metadata& table_meta) if (is_fixed_width(column.type())) { return size_of(column.type()) * column.size(); - } else if (column.type().id() == type_id::STRING) { - auto const scol = strings_column_view(column); + } else if (column.type().id() == type_id::STRING or column.type().id() == type_id::BINARY) { + auto const offsets = column.type().id() == type_id::STRING + ? strings_column_view(column).offsets() + : binary_column_view(column).offsets(); return cudf::strings::detail::get_offset_value( - scol.offsets(), column.size() + column.offset(), stream) - - cudf::strings::detail::get_offset_value(scol.offsets(), column.offset(), stream); + offsets, column.size() + column.offset(), stream) - + cudf::strings::detail::get_offset_value(offsets, column.offset(), stream); } else if (column.type().id() == type_id::STRUCT) { auto const scol = structs_column_view(column); size_t ret = 0; diff --git a/cpp/src/io/statistics/statistics_type_identification.cuh b/cpp/src/io/statistics/statistics_type_identification.cuh index 4324e9eb4ec8..1deaf951a07a 100644 --- a/cpp/src/io/statistics/statistics_type_identification.cuh +++ b/cpp/src/io/statistics/statistics_type_identification.cuh @@ -13,6 +13,7 @@ #include "byte_array_view.cuh" #include "conversion_type_select.cuh" +#include #include #include #include @@ -118,7 +119,9 @@ class extrema_type { typename std::conditional_t< std::is_same_v, string_view, - std::conditional_t, byte_array_view, void>>>; + std::conditional_t or std::is_same_v, + byte_array_view, + void>>>; // unsigned int/bool -> uint64_t // signed int -> int64_t @@ -133,7 +136,8 @@ class extrema_type { // Does type T have an extrema? static constexpr bool is_supported = std::is_arithmetic_v or std::is_same_v or cudf::is_duration() or - cudf::is_timestamp() or cudf::is_fixed_point() or std::is_same_v; + cudf::is_timestamp() or cudf::is_fixed_point() or std::is_same_v or + std::is_same_v; using type = typename std:: conditional_t, arithmetic_extrema_type, non_arithmetic_extrema_type>; @@ -146,6 +150,9 @@ class extrema_type { if constexpr (std::is_arithmetic_v or std::is_same_v or std::is_same_v) { return val; + } else if constexpr (std::is_same_v) { + return byte_array_view{reinterpret_cast(val.data()), + static_cast(val.size_bytes())}; } else if constexpr (cudf::is_fixed_point()) { return val.value(); } else if constexpr (cudf::is_duration()) { @@ -176,7 +183,8 @@ class aggregation_type { using non_arithmetic_aggregation_type = typename std::conditional_t< cudf::is_fixed_point() or cudf::is_duration() or cudf::is_timestamp() // To be disabled with static_assert - or std::is_same_v or std::is_same_v, + or std::is_same_v or std::is_same_v or + std::is_same_v, typename std::conditional_t, __int128_t, int64_t>, void>; @@ -194,7 +202,8 @@ class aggregation_type { // Does type T aggregate? static constexpr bool is_supported = std::is_arithmetic_v or std::is_same_v or cudf::is_duration() or cudf::is_fixed_point() or - std::is_same_v; + std::is_same_v or + std::is_same_v; using type = typename std::conditional_t, arithmetic_aggregation_type, @@ -205,7 +214,8 @@ class aggregation_type { */ __device__ static type convert(T const& val) { - if constexpr (std::is_same_v or std::is_same_v) { + if constexpr (std::is_same_v or std::is_same_v or + std::is_same_v) { return val.size_bytes(); } else if constexpr (std::is_integral_v) { return val; diff --git a/cpp/src/io/utilities/column_buffer.cpp b/cpp/src/io/utilities/column_buffer.cpp index c1af6be6acd4..d56af5f903d2 100644 --- a/cpp/src/io/utilities/column_buffer.cpp +++ b/cpp/src/io/utilities/column_buffer.cpp @@ -10,6 +10,7 @@ #include "column_buffer.hpp" +#include #include #include #include @@ -66,7 +67,7 @@ void cudf::io::detail::inline_column_buffer::create_string_data(size_t num_bytes namespace { /** - * @brief Recursively copy `name`, `user_data`, and `string_as_binary` fields of one buffer to + * @brief Recursively copy metadata fields of one buffer to * another. * * @param buff The old output buffer @@ -75,9 +76,9 @@ namespace { template void copy_buffer_data(string_policy const& buff, string_policy& new_buff) { - new_buff.name = buff.name; - new_buff.user_data = buff.user_data; - new_buff.string_as_binary = buff.string_as_binary; + new_buff.name = buff.name; + new_buff.user_data = buff.user_data; + new_buff.byte_array_output = buff.byte_array_output; for (auto const& child : buff.children) { auto& new_child = new_buff.children.emplace_back(string_policy(child.type, child.is_nullable)); copy_buffer_data(child, new_child); @@ -190,50 +191,64 @@ std::unique_ptr make_column(column_buffer_base& buffer, } switch (buffer.type.id()) { case type_id::STRING: { - if (schema.value_or(reader_column_schema{}).is_enabled_convert_binary_to_strings() and - not buffer.string_as_binary) { + auto const schema_output = schema.has_value() ? schema->get_byte_array_output() + : byte_array_output_type::DEFAULT; + auto const requested_output = schema_output == byte_array_output_type::DEFAULT + ? buffer.byte_array_output + : schema_output; + auto const output = requested_output == byte_array_output_type::DEFAULT + ? byte_array_output_type::STRING + : requested_output; + if (output == byte_array_output_type::STRING) { if (schema_info != nullptr) { schema_info->children.emplace_back("offsets"); } - // make_strings_column allocates new memory, it does not simply move - // from the inputs, so we need to pass it the memory resource given to - // the buffer on construction so that the memory is allocated using the - // resource that the calling code expected. return buffer.make_string_column(stream); - } else { - // convert to binary - auto const string_col = buffer.make_string_column(stream); - auto const num_rows = string_col->size(); - auto const null_count = string_col->null_count(); - auto col_content = string_col->release(); - - // convert to uint8 column, strings are currently stored as int8 - auto data = col_content.data.release(); - auto char_size = data->size(); - - CUDF_EXPECTS(char_size < static_cast(std::numeric_limits::max()), - "Cannot convert strings column to lists column due to size_type limit", - std::overflow_error); - - auto uint8_col = std::make_unique( - data_type{type_id::UINT8}, char_size, std::move(*data), rmm::device_buffer{}, 0); + } - if (schema_info != nullptr) { - schema_info->children.emplace_back("offsets"); - schema_info->children.emplace_back("binary"); - // cuDF type will be list, but remember it was originally binary data - schema_info->is_binary = true; - if (schema.has_value() and schema->get_type_length() > 0) { - schema_info->type_length = schema->get_type_length(); - } - } + auto const string_col = buffer.make_string_column(stream); + auto const num_rows = string_col->size(); + auto const null_count = string_col->null_count(); + auto col_content = string_col->release(); - return make_lists_column( + if (output == byte_array_output_type::BINARY) { + if (schema_info != nullptr) { schema_info->children.emplace_back("offsets"); } + return make_binary_column( num_rows, std::move(col_content.children[strings_column_view::offsets_column_index]), - std::move(uint8_col), + std::move(*col_content.data), null_count, std::move(*col_content.null_mask)); } + + CUDF_EXPECTS(output == byte_array_output_type::LIST_UINT8, + "Invalid byte-array output type"); + // Convert to uint8 column; strings are currently stored as int8. + auto data = col_content.data.release(); + auto char_size = data->size(); + + CUDF_EXPECTS(char_size < static_cast(std::numeric_limits::max()), + "Cannot convert strings column to lists column due to size_type limit", + std::overflow_error); + + auto uint8_col = std::make_unique( + data_type{type_id::UINT8}, char_size, std::move(*data), rmm::device_buffer{}, 0); + + if (schema_info != nullptr) { + schema_info->children.emplace_back("offsets"); + schema_info->children.emplace_back("binary"); + // cuDF type will be list, but remember it was originally binary data. + schema_info->is_binary = true; + if (schema.has_value() and schema->get_type_length() > 0) { + schema_info->type_length = schema->get_type_length(); + } + } + + return make_lists_column( + num_rows, + std::move(col_content.children[strings_column_view::offsets_column_index]), + std::move(uint8_col), + null_count, + std::move(*col_content.null_mask)); } case type_id::LIST: { @@ -321,7 +336,11 @@ std::unique_ptr empty_like(column_buffer_base& buffer, switch (buffer.type.id()) { case type_id::STRING: { - if (buffer.string_as_binary) { + if (buffer.byte_array_output == byte_array_output_type::BINARY) { + if (schema_info != nullptr) { schema_info->children.emplace_back("offsets"); } + return make_empty_binary_column(); + } + if (buffer.byte_array_output == byte_array_output_type::LIST_UINT8) { auto offsets = cudf::make_empty_column(type_id::INT32); auto child = cudf::make_empty_column(type_id::UINT8); if (schema_info != nullptr) { diff --git a/cpp/src/io/utilities/column_buffer.hpp b/cpp/src/io/utilities/column_buffer.hpp index 1bc53a3a39e5..8beacc6b8810 100644 --- a/cpp/src/io/utilities/column_buffer.hpp +++ b/cpp/src/io/utilities/column_buffer.hpp @@ -161,8 +161,8 @@ class column_buffer_base { bool is_nullable{false}; size_type size{0}; uint32_t user_data{0}; // arbitrary user data - // Materialize a STRING/BYTE_ARRAY-backed buffer as a `list` column - bool string_as_binary{false}; + // Output representation for a STRING-layout Parquet byte-array buffer. + byte_array_output_type byte_array_output{byte_array_output_type::DEFAULT}; std::string name; std::vector children; diff --git a/cpp/src/reductions/minmax.cu b/cpp/src/reductions/minmax.cu index f643ce6aa11a..fe7e01b23ccf 100644 --- a/cpp/src/reductions/minmax.cu +++ b/cpp/src/reductions/minmax.cu @@ -178,7 +178,7 @@ struct minmax_dictionary_functor { static constexpr bool is_supported() { return !cudf::is_dictionary() && !std::is_same_v && - !std::is_same_v; + !std::is_same_v && !std::is_same_v; } template @@ -229,7 +229,8 @@ struct minmax_functor { template static constexpr bool is_supported() { - return !(std::is_same_v || std::is_same_v); + return !(std::is_same_v || std::is_same_v || + std::is_same_v); } template diff --git a/cpp/src/reductions/scan/scan_inclusive.cu b/cpp/src/reductions/scan/scan_inclusive.cu index 680763ecbf68..c0e0170078ff 100644 --- a/cpp/src/reductions/scan/scan_inclusive.cu +++ b/cpp/src/reductions/scan/scan_inclusive.cu @@ -159,7 +159,8 @@ struct scan_dispatcher { if constexpr (std::is_same_v) { return std::is_same_v || std::is_same_v; } else { - return std::is_invocable_v && !cudf::is_dictionary(); + return std::is_invocable_v && !cudf::is_dictionary() && + !std::is_same_v; } } diff --git a/cpp/src/reductions/segmented/simple.cuh b/cpp/src/reductions/segmented/simple.cuh index 3668c3b53d94..a793150158d0 100644 --- a/cpp/src/reductions/segmented/simple.cuh +++ b/cpp/src/reductions/segmented/simple.cuh @@ -321,7 +321,8 @@ struct same_column_type_dispatcher { static constexpr bool is_supported() { return !(cudf::is_dictionary() || std::is_same_v || - std::is_same_v); + std::is_same_v || + std::is_same_v); } public: diff --git a/cpp/src/reductions/simple.cuh b/cpp/src/reductions/simple.cuh index 7ea6b2b608a9..81f8bd3507fd 100644 --- a/cpp/src/reductions/simple.cuh +++ b/cpp/src/reductions/simple.cuh @@ -271,7 +271,8 @@ struct same_element_type_dispatcher { template static constexpr bool is_supported() { - return !cudf::is_dictionary() && !std::is_same_v; + return !cudf::is_dictionary() && + !std::is_same_v && !std::is_same_v; } template diff --git a/cpp/src/replace/clamp.cu b/cpp/src/replace/clamp.cu index 79986de1c11e..168636a527e4 100644 --- a/cpp/src/replace/clamp.cu +++ b/cpp/src/replace/clamp.cu @@ -303,6 +303,19 @@ std::unique_ptr dispatch_clamp::operator()( CUDF_FAIL("clamp for list_view not supported"); } +template <> +std::unique_ptr dispatch_clamp::operator()( + column_view const&, + scalar const&, + scalar const&, + scalar const&, + scalar const&, + rmm::cuda_stream_view, + rmm::device_async_resource_ref) +{ + CUDF_FAIL("clamp for BINARY is not yet supported"); +} + template <> std::unique_ptr dispatch_clamp::operator()(column_view const& input, scalar const& lo, diff --git a/cpp/src/rolling/detail/rolling_operators.cuh b/cpp/src/rolling/detail/rolling_operators.cuh index f3d39f1edb06..1892a9d6c21b 100644 --- a/cpp/src/rolling/detail/rolling_operators.cuh +++ b/cpp/src/rolling/detail/rolling_operators.cuh @@ -122,8 +122,9 @@ struct DeviceRollingArgMinMaxBase { static constexpr bool is_supported() { // Right now only support ARGMIN/ARGMAX for compound-types but not lists - auto const type_supported = - cudf::is_compound() && !std::is_same_v; + auto const type_supported = cudf::is_compound() && + !std::is_same_v && + !std::is_same_v; auto const op_supported = op == aggregation::Kind::ARGMIN || op == aggregation::Kind::ARGMAX; return type_supported && op_supported; @@ -238,7 +239,7 @@ struct DeviceRollingArgMinMaxDictionary : DeviceRollingArgMinMaxBase - requires(cudf::is_dictionary_key()) + requires(cudf::is_dictionary_key() && !std::is_same_v) size_type __device__ operator()(column_device_view const& dict, bool has_nulls, size_type start_index, @@ -263,7 +264,7 @@ struct DeviceRollingArgMinMaxDictionary : DeviceRollingArgMinMaxBase= min_periods ? index : -1; } template - requires(not cudf::is_dictionary_key()) + requires(not cudf::is_dictionary_key() || std::is_same_v) size_type __device__ operator()(column_device_view const&, bool, size_type, size_type, size_type) { diff --git a/cpp/src/search/contains_scalar.cu b/cpp/src/search/contains_scalar.cu index c6f329e5e1f8..0a0fbfdf124a 100644 --- a/cpp/src/search/contains_scalar.cu +++ b/cpp/src/search/contains_scalar.cu @@ -51,9 +51,8 @@ struct contains_scalar_dispatch { // SFINAE with conditional return type because we need to support device lambda in this function. // This is required due to a limitation of nvcc. template - std::enable_if_t(), bool> operator()(column_view const& haystack, - scalar const& needle, - cuda::stream_ref stream) const + std::enable_if_t() && !std::is_same_v, bool> operator()( + column_view const& haystack, scalar const& needle, cuda::stream_ref stream) const { CUDF_EXPECTS(cudf::have_same_types(haystack, needle), "Scalar and column types must match", @@ -80,6 +79,13 @@ struct contains_scalar_dispatch { stream) > 0; } + template + std::enable_if_t, bool> operator()( + column_view const&, scalar const&, cuda::stream_ref) const + { + CUDF_FAIL("BINARY scalar support is not yet implemented"); + } + template std::enable_if_t(), bool> operator()(column_view const& haystack, scalar const& needle, diff --git a/cpp/src/utilities/traits.cpp b/cpp/src/utilities/traits.cpp index 389ea9b69779..9b57914e8c38 100644 --- a/cpp/src/utilities/traits.cpp +++ b/cpp/src/utilities/traits.cpp @@ -3,6 +3,7 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include #include #include #include diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 06c0462ebed9..d6c5581157e3 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -95,6 +95,10 @@ include(../cmake/thirdparty/get_arrow.cmake) # test sources ################################################################################## # ################################################################################################## +# ################################################################################################## +# * binary tests ---------------------------------------------------------------------------------- +ConfigureTest(BINARY_TEST binary/binary_test.cu) + # ################################################################################################## # * column tests ---------------------------------------------------------------------------------- ConfigureTest( diff --git a/cpp/tests/binary/binary_test.cu b/cpp/tests/binary/binary_test.cu new file mode 100644 index 000000000000..8ef54693fd27 --- /dev/null +++ b/cpp/tests/binary/binary_test.cu @@ -0,0 +1,361 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +#include + +#include +#include +#include +#include + +namespace { + +struct BinaryTest : public cudf::test::BaseFixture {}; + +struct is_binary_dispatch { + template + bool operator()() const + { + return std::is_same_v; + } +}; + +rmm::device_buffer make_payload(std::vector const& bytes, + rmm::cuda_stream_view stream) +{ + rmm::device_buffer result(bytes.size(), stream); + if (not bytes.empty()) { + CUDF_CUDA_TRY(cudaMemcpyAsync( + result.data(), bytes.data(), bytes.size(), cudaMemcpyHostToDevice, stream.value())); + } + return result; +} + +template +std::unique_ptr make_test_column(rmm::cuda_stream_view stream) +{ + auto offsets = cudf::test::fixed_width_column_wrapper({0, 2, 2, 5}).release(); + auto payload = make_payload({0x00, 0xFF, 0x61, 0x62, 0x63}, stream); + return cudf::make_binary_column( + 3, std::move(offsets), std::move(payload), 0, rmm::device_buffer{}); +} + +CUDF_KERNEL void inspect_binary(cudf::column_device_view const* input, + cudf::size_type* sizes, + uint8_t* first_bytes) +{ + auto const index = static_cast(threadIdx.x); + if (index < input->size()) { + auto const value = input->element(index); + sizes[index] = value.size_bytes(); + first_bytes[index] = value.empty() ? uint8_t{0} : value[0]; + } +} + +TEST_F(BinaryTest, BinaryViewUsesUnsignedLexicographicOrdering) +{ + std::array low{0x00, 0xFF}; + std::array high{0x80, 0x00}; + std::array extension{0x00, 0xFF, 0x00}; + + auto const low_view = cudf::binary_view{low.data(), low.size()}; + auto const high_view = cudf::binary_view{high.data(), high.size()}; + auto const extension_view = cudf::binary_view{extension.data(), extension.size()}; + + EXPECT_LT(low_view, high_view); + EXPECT_LT(low_view, extension_view); + EXPECT_EQ(low_view, (cudf::binary_view{low.data(), low.size()})); +} + +TEST_F(BinaryTest, TypeDispatchesToBinaryView) +{ + static_assert(cudf::type_to_id() == cudf::type_id::BINARY); + EXPECT_TRUE( + cudf::type_dispatcher(cudf::data_type{cudf::type_id::BINARY}, is_binary_dispatch{})); + EXPECT_EQ(cudf::type_to_name(cudf::data_type{cudf::type_id::BINARY}), "cudf::binary_view"); +} + +TEST_F(BinaryTest, TypeTraitsDescribeVariableWidthLeaf) +{ + auto const type = cudf::data_type{cudf::type_id::BINARY}; + + EXPECT_TRUE(cudf::is_compound(type)); + EXPECT_FALSE(cudf::is_nested(type)); + EXPECT_FALSE(cudf::is_fixed_width(type)); + EXPECT_TRUE(cudf::is_equality_comparable(type)); + EXPECT_TRUE(cudf::is_relationally_comparable(type)); +} + +TEST_F(BinaryTest, FactoryAcceptsInt32Offsets) +{ + auto const stream = cudf::get_default_stream(); + auto column = make_test_column(stream); + auto view = cudf::binary_column_view{column->view()}; + + EXPECT_EQ(column->type().id(), cudf::type_id::BINARY); + EXPECT_EQ(column->size(), 3); + EXPECT_EQ(column->num_children(), 1); + EXPECT_EQ(view.offsets().type().id(), cudf::type_id::INT32); + EXPECT_EQ(view.bytes_size(stream), 5); +} + +TEST_F(BinaryTest, FactoryAcceptsInt64Offsets) +{ + auto const stream = cudf::get_default_stream(); + auto column = make_test_column(stream); + auto view = cudf::binary_column_view{column->view()}; + + EXPECT_EQ(view.offsets().type().id(), cudf::type_id::INT64); + EXPECT_EQ(view.bytes_size(stream), 5); +} + +TEST_F(BinaryTest, EmptyColumnHasCanonicalLayout) +{ + auto column = cudf::make_empty_binary_column(); + + EXPECT_EQ(column->type().id(), cudf::type_id::BINARY); + EXPECT_EQ(column->size(), 0); + EXPECT_EQ(column->num_children(), 0); + EXPECT_EQ(cudf::binary_column_view{column->view()}.bytes_size(cudf::get_default_stream()), 0); +} + +TEST_F(BinaryTest, FactoryRejectsInvalidOffsetsType) +{ + auto offsets = cudf::test::fixed_width_column_wrapper({0, 1}).release(); + auto payload = make_payload({0x01}, cudf::get_default_stream()); + + EXPECT_THROW(cudf::make_binary_column( + 1, std::move(offsets), std::move(payload), 0, rmm::device_buffer{}), + cudf::logic_error); +} + +TEST_F(BinaryTest, DeviceViewAccessesRows) +{ + auto const stream = cudf::get_default_stream(); + auto column = make_test_column(stream); + auto device_view = cudf::column_device_view::create(column->view(), stream); + rmm::device_uvector sizes(column->size(), stream); + rmm::device_uvector first_bytes(column->size(), stream); + + inspect_binary<<<1, column->size(), 0, stream.value()>>>( + device_view.get(), sizes.data(), first_bytes.data()); + CUDF_CUDA_TRY(cudaGetLastError()); + + auto const host_sizes = cudf::detail::make_std_vector_async(sizes, stream); + auto const host_first = cudf::detail::make_std_vector_async(first_bytes, stream); + stream.synchronize(); + + EXPECT_EQ(host_sizes, (std::vector{2, 0, 3})); + EXPECT_EQ(host_first, (std::vector{0x00, 0x00, 0x61})); +} + +TEST_F(BinaryTest, OwningCopyNormalizesSlicedOffsets) +{ + auto const stream = cudf::get_default_stream(); + auto input = make_test_column(stream); + auto const sliced = cudf::slice(input->view(), {1, 3}, stream).front(); + auto copied = std::make_unique(sliced, stream); + auto copied_view = cudf::binary_column_view{copied->view()}; + + EXPECT_EQ(copied->size(), 2); + EXPECT_EQ(copied_view.bytes_size(stream), 3); + EXPECT_EQ(cudf::detail::get_value(copied_view.offsets(), 0, stream), 0); + EXPECT_EQ(cudf::detail::get_value(copied_view.offsets(), 2, stream), 3); + + auto device_view = cudf::column_device_view::create(copied->view(), stream); + rmm::device_uvector sizes(copied->size(), stream); + rmm::device_uvector first_bytes(copied->size(), stream); + inspect_binary<<<1, copied->size(), 0, stream.value()>>>( + device_view.get(), sizes.data(), first_bytes.data()); + CUDF_CUDA_TRY(cudaGetLastError()); + + auto const host_sizes = cudf::detail::make_std_vector_async(sizes, stream); + auto const host_first = cudf::detail::make_std_vector_async(first_bytes, stream); + stream.synchronize(); + + EXPECT_EQ(host_sizes, (std::vector{0, 3})); + EXPECT_EQ(host_first, (std::vector{0x00, 0x61})); +} + +TEST_F(BinaryTest, GatherReordersBinaryRows) +{ + auto const stream = cudf::get_default_stream(); + auto input = make_test_column(stream); + auto gather_map = cudf::test::fixed_width_column_wrapper({2, 0, 1}); + auto result = cudf::gather(cudf::table_view{{input->view()}}, + gather_map, + cudf::out_of_bounds_policy::DONT_CHECK, + stream); + auto output = std::move(result->release().front()); + + auto device_view = cudf::column_device_view::create(output->view(), stream); + rmm::device_uvector sizes(output->size(), stream); + rmm::device_uvector first_bytes(output->size(), stream); + inspect_binary<<<1, output->size(), 0, stream.value()>>>( + device_view.get(), sizes.data(), first_bytes.data()); + CUDF_CUDA_TRY(cudaGetLastError()); + + auto const host_sizes = cudf::detail::make_std_vector_async(sizes, stream); + auto const host_first = cudf::detail::make_std_vector_async(first_bytes, stream); + stream.synchronize(); + + EXPECT_EQ(host_sizes, (std::vector{3, 2, 0})); + EXPECT_EQ(host_first, (std::vector{0x61, 0x00, 0x00})); +} + +TEST_F(BinaryTest, GatherNullifiesOutOfBoundsRows) +{ + auto const stream = cudf::get_default_stream(); + auto input = make_test_column(stream); + auto gather_map = cudf::test::fixed_width_column_wrapper({0, 99, 2}); + auto result = cudf::gather(cudf::table_view{{input->view()}}, + gather_map, + cudf::out_of_bounds_policy::NULLIFY, + stream); + auto output = std::move(result->release().front()); + + EXPECT_EQ(output->null_count(), 1); + EXPECT_EQ(cudf::binary_column_view{output->view()}.bytes_size(stream), 5); +} + +TEST_F(BinaryTest, ConcatenateSupportsSlicedColumns) +{ + auto const stream = cudf::get_default_stream(); + auto input = make_test_column(stream); + auto const sliced = cudf::slice(input->view(), {1, 3}, stream).front(); + auto output = + cudf::concatenate(std::vector{sliced, input->view()}, stream); + + auto device_view = cudf::column_device_view::create(output->view(), stream); + rmm::device_uvector sizes(output->size(), stream); + rmm::device_uvector first_bytes(output->size(), stream); + inspect_binary<<<1, output->size(), 0, stream.value()>>>( + device_view.get(), sizes.data(), first_bytes.data()); + CUDF_CUDA_TRY(cudaGetLastError()); + + auto const host_sizes = cudf::detail::make_std_vector_async(sizes, stream); + stream.synchronize(); + + EXPECT_EQ(host_sizes, (std::vector{0, 3, 2, 0, 3})); + EXPECT_EQ(cudf::binary_column_view{output->view()}.bytes_size(stream), 8); +} + +TEST_F(BinaryTest, ScatterMovesBinaryRows) +{ + auto const stream = cudf::get_default_stream(); + auto target = make_test_column(stream); + auto const source = cudf::slice(target->view(), {2, 3}, stream).front(); + auto scatter_map = cudf::test::fixed_width_column_wrapper({1}); + auto result = cudf::scatter(cudf::table_view{{source}}, + scatter_map, + cudf::table_view{{target->view()}}, + stream); + auto output = std::move(result->release().front()); + + auto device_view = cudf::column_device_view::create(output->view(), stream); + rmm::device_uvector sizes(output->size(), stream); + rmm::device_uvector first_bytes(output->size(), stream); + inspect_binary<<<1, output->size(), 0, stream.value()>>>( + device_view.get(), sizes.data(), first_bytes.data()); + CUDF_CUDA_TRY(cudaGetLastError()); + + auto const host_sizes = cudf::detail::make_std_vector_async(sizes, stream); + stream.synchronize(); + + EXPECT_EQ(host_sizes, (std::vector{2, 3, 3})); + EXPECT_EQ(cudf::binary_column_view{output->view()}.bytes_size(stream), 8); +} + +TEST_F(BinaryTest, HashesBinaryPayloadBytes) +{ + auto const stream = cudf::get_default_stream(); + auto offsets = cudf::test::fixed_width_column_wrapper({0, 2, 4, 4}).release(); + auto payload = make_payload({0x00, 0xFF, 0x00, 0xFF}, stream); + auto input = cudf::make_binary_column( + 3, std::move(offsets), std::move(payload), 0, rmm::device_buffer{}); + + auto hashes = cudf::hashing::murmurhash3_x86_32(cudf::table_view{{input->view()}}, 0, stream); + auto const [host_hashes, validity] = cudf::test::to_host(hashes->view(), stream); + + EXPECT_EQ(host_hashes[0], host_hashes[1]); + EXPECT_NE(host_hashes[0], host_hashes[2]); +} + +TEST_F(BinaryTest, SortsByUnsignedByteOrder) +{ + auto const stream = cudf::get_default_stream(); + auto input = make_test_column(stream); + + auto order = cudf::sorted_order(cudf::table_view{{input->view()}}, {}, {}, stream); + auto const [host_order, validity] = + cudf::test::to_host(order->view(), stream); + + EXPECT_EQ(host_order, (std::vector{1, 0, 2})); +} + +TEST_F(BinaryTest, DistinctUsesBinaryEquality) +{ + auto const stream = cudf::get_default_stream(); + auto offsets = cudf::test::fixed_width_column_wrapper({0, 2, 4, 4}).release(); + auto payload = make_payload({0x00, 0xFF, 0x00, 0xFF}, stream); + auto input = cudf::make_binary_column( + 3, std::move(offsets), std::move(payload), 0, rmm::device_buffer{}); + + auto result = cudf::distinct(cudf::table_view{{input->view()}}, + {0}, + cudf::duplicate_keep_option::KEEP_FIRST, + cudf::null_equality::EQUAL, + cudf::nan_equality::ALL_EQUAL, + stream); + + EXPECT_EQ(result->num_rows(), 2); + EXPECT_EQ(result->get_column(0).type().id(), cudf::type_id::BINARY); +} + +TEST_F(BinaryTest, DictionaryRoundTripPreservesBinary) +{ + auto const stream = cudf::get_default_stream(); + auto offsets = cudf::test::fixed_width_column_wrapper({0, 2, 4, 4}).release(); + auto payload = make_payload({0x00, 0xFF, 0x00, 0xFF}, stream); + auto input = cudf::make_binary_column( + 3, std::move(offsets), std::move(payload), 0, rmm::device_buffer{}); + + auto encoded = + cudf::dictionary::encode(input->view(), cudf::data_type{cudf::type_id::INT32}, stream); + auto decoded = cudf::dictionary::decode(cudf::dictionary_column_view{encoded->view()}, stream); + + EXPECT_EQ(decoded->type().id(), cudf::type_id::BINARY); + EXPECT_EQ(decoded->size(), input->size()); + EXPECT_EQ(cudf::binary_column_view{decoded->view()}.bytes_size(stream), 4); +} + +} // namespace + +CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/io/parquet_reader_test.cpp b/cpp/tests/io/parquet_reader_test.cpp index 08e48dd213eb..48afe9f0e1f9 100644 --- a/cpp/tests/io/parquet_reader_test.cpp +++ b/cpp/tests/io/parquet_reader_test.cpp @@ -1593,7 +1593,14 @@ TEST_F(ParquetReaderTest, BinaryAsStrings) cudf::io::parquet_reader_options in_opts = cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}) - .set_column_schema({{}, {}, {}, {}, {}}); + .set_column_schema( + {{}, + {}, + {}, + cudf::io::reader_column_schema().set_byte_array_output( + cudf::io::byte_array_output_type::STRING), + cudf::io::reader_column_schema().set_byte_array_output( + cudf::io::byte_array_output_type::STRING)}); auto result = cudf::io::read_parquet(in_opts); CUDF_TEST_EXPECT_TABLES_EQUAL(expected_string, result.tbl->view()); @@ -1602,7 +1609,33 @@ TEST_F(ParquetReaderTest, BinaryAsStrings) cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}); result = cudf::io::read_parquet(default_in_opts); - CUDF_TEST_EXPECT_TABLES_EQUAL(expected_string, result.tbl->view()); + EXPECT_EQ(result.tbl->view().column(1).type().id(), cudf::type_id::STRING); + EXPECT_EQ(result.tbl->view().column(3).type().id(), cudf::type_id::BINARY); + EXPECT_EQ(result.tbl->view().column(4).type().id(), cudf::type_id::BINARY); + + auto native_binary_filepath = temp_env->get_temp_filepath("NativeBinaryWrite.parquet"); + cudf::io::write_parquet(cudf::io::parquet_writer_options::builder( + cudf::io::sink_info{native_binary_filepath}, result.tbl->view()) + .write_arrow_schema(true)); + auto native_binary_result = cudf::io::read_parquet( + cudf::io::parquet_reader_options::builder(cudf::io::source_info{native_binary_filepath}) + .use_arrow_schema(true)); + CUDF_TEST_EXPECT_TABLES_EQUAL(result.tbl->view(), native_binary_result.tbl->view()); + EXPECT_EQ(native_binary_result.tbl->view().column(3).type().id(), cudf::type_id::BINARY); + EXPECT_EQ(native_binary_result.tbl->view().column(4).type().id(), cudf::type_id::BINARY); + + auto binary_override = cudf::io::parquet_reader_options::builder( + cudf::io::source_info{filepath}) + .set_column_schema( + {{}, + cudf::io::reader_column_schema().set_byte_array_output( + cudf::io::byte_array_output_type::BINARY), + {}, + {}, + {}}) + .build(); + result = cudf::io::read_parquet(binary_override); + EXPECT_EQ(result.tbl->view().column(1).type().id(), cudf::type_id::BINARY); std::vector md{ {}, @@ -1682,6 +1715,10 @@ TEST_F(ParquetReaderTest, NestedByteArray) auto result = cudf::io::read_parquet(in_opts); CUDF_TEST_EXPECT_TABLES_EQUAL(expected, result.tbl->view()); + + auto default_result = cudf::io::read_parquet( + cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath})); + EXPECT_EQ(default_result.tbl->view().column(2).child(1).type().id(), cudf::type_id::BINARY); } TEST_F(ParquetReaderTest, StructByteArray) @@ -1720,6 +1757,10 @@ TEST_F(ParquetReaderTest, StructByteArray) auto result = cudf::io::read_parquet(in_opts); CUDF_TEST_EXPECT_TABLES_EQUAL(expected, result.tbl->view()); + + auto default_result = cudf::io::read_parquet( + cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath})); + EXPECT_EQ(default_result.tbl->view().column(0).child(0).type().id(), cudf::type_id::BINARY); } TEST_F(ParquetReaderTest, NestingOptimizationTest) diff --git a/cpp/tests/io/parquet_writer_test.cpp b/cpp/tests/io/parquet_writer_test.cpp index ec0267c3bfea..f571a7578a16 100644 --- a/cpp/tests/io/parquet_writer_test.cpp +++ b/cpp/tests/io/parquet_writer_test.cpp @@ -2389,11 +2389,33 @@ TEST_F(ParquetWriterTest, WriteFixedLenByteArray) cudf::io::write_parquet(out_opts); cudf::io::parquet_reader_options in_opts = - cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}); + cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}) + .set_column_schema( + std::vector( + 4, cudf::io::reader_column_schema().set_convert_binary_to_strings(false))); auto result = cudf::io::read_parquet(in_opts); CUDF_TEST_EXPECT_TABLES_EQUAL(expected, result.tbl->view()); + auto binary_result = cudf::io::read_parquet( + cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath})); + for (auto const& column : binary_result.tbl->view()) { + EXPECT_EQ(column.type().id(), cudf::type_id::BINARY); + } + + cudf::io::table_input_metadata binary_metadata(binary_result.tbl->view()); + for (auto& column : binary_metadata.column_metadata) { + column.set_type_length(fixed_width); + } + auto binary_filepath = temp_env->get_temp_filepath("WriteNativeBinaryFixedLenByteArray.parquet"); + cudf::io::write_parquet( + cudf::io::parquet_writer_options::builder( + cudf::io::sink_info{binary_filepath}, binary_result.tbl->view()) + .metadata(std::move(binary_metadata))); + auto binary_roundtrip = cudf::io::read_parquet( + cudf::io::parquet_reader_options::builder(cudf::io::source_info{binary_filepath})); + CUDF_TEST_EXPECT_TABLES_EQUAL(binary_result.tbl->view(), binary_roundtrip.tbl->view()); + // check page headers to make sure each column is encoded with the appropriate encoder auto const source = cudf::io::datasource::create(filepath); cudf::io::parquet::FileMetaData fmd; diff --git a/cpp/tests/utilities/debug_utilities.cu b/cpp/tests/utilities/debug_utilities.cu index 9241ce25a30e..97690d88a2c9 100644 --- a/cpp/tests/utilities/debug_utilities.cu +++ b/cpp/tests/utilities/debug_utilities.cu @@ -209,6 +209,26 @@ struct column_view_printer { } } + template + void operator()(cudf::column_view const& col, + std::vector& out, + std::string const&, + rmm::cuda_stream_view stream, + cudf::memory_resources mr) + requires(std::is_same_v) + { + auto const host_mask = cudf::test::bitmask_to_host(col, stream, mr); + out.resize(col.size()); + std::transform(cuda::counting_iterator{0}, + cuda::counting_iterator{col.size()}, + out.begin(), + [&](auto idx) { + return host_mask.empty() || bit_is_set(host_mask.data(), idx) + ? std::string{""} + : std::string{"NULL"}; + }); + } + template void operator()(cudf::column_view const& col, std::vector& out,