diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt index 192ed1b978b6..365d073faf86 100644 --- a/cpp/benchmarks/CMakeLists.txt +++ b/cpp/benchmarks/CMakeLists.txt @@ -176,6 +176,10 @@ ConfigureNVBench( SEARCH_NVBENCH search/contains_scalar.cpp search/contains_table.cpp search/search.cpp ) +# ################################################################################################## +# * scalar benchmark ------------------------------------------------------------------------------ +ConfigureNVBench(SCALAR_NVBENCH scalar/scalar.cpp) + # ################################################################################################## # * sort benchmark -------------------------------------------------------------------------------- ConfigureNVBench( diff --git a/cpp/benchmarks/scalar/scalar.cpp b/cpp/benchmarks/scalar/scalar.cpp new file mode 100644 index 000000000000..37b6eb9ce5e9 --- /dev/null +++ b/cpp/benchmarks/scalar/scalar.cpp @@ -0,0 +1,55 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include +#include + +#include + +#include +#include + +namespace { + +void numeric_scalar_construction(nvbench::state& state) +{ + auto const is_valid = static_cast(state.get_int64("valid")); + + auto const mem_stats_logger = cudf::memory_stats_logger(); + state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { + auto const stream = rmm::cuda_stream_view{launch.get_stream()}; + [[maybe_unused]] auto scalar = cudf::numeric_scalar{ + 42, is_valid, stream, cudf::get_current_device_resource_ref()}; + }); + state.add_buffer_size( + mem_stats_logger.peak_memory_usage(), "peak_memory_usage", "peak_memory_usage"); +} + +void string_scalar_construction(nvbench::state& state) +{ + auto const num_bytes = state.get_int64("num_bytes"); + auto const value = std::string(static_cast(num_bytes), 'a'); + + auto const mem_stats_logger = cudf::memory_stats_logger(); + state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { + auto const stream = rmm::cuda_stream_view{launch.get_stream()}; + [[maybe_unused]] auto scalar = cudf::string_scalar{ + value, true, stream, cudf::get_current_device_resource_ref()}; + }); + state.add_buffer_size( + mem_stats_logger.peak_memory_usage(), "peak_memory_usage", "peak_memory_usage"); +} + +} // namespace + +NVBENCH_BENCH(numeric_scalar_construction) + .set_name("numeric_scalar_construction") + .add_int64_axis("valid", {0, 1}); + +NVBENCH_BENCH(string_scalar_construction) + .set_name("string_scalar_construction") + .add_int64_power_of_two_axis("num_bytes", {0, 4, 8, 12, 16, 20}); diff --git a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md index 6624b54038fc..db3d242928d7 100644 --- a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md +++ b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md @@ -403,22 +403,23 @@ Use `cudf::host_span` only when one of the following applies: A `cudf::scalar` is an object that can represent a singular, nullable value of any of the types currently supported by cudf. Each type of value is represented by a separate type of scalar class which are all derived from `cudf::scalar`. e.g. A `numeric_scalar` holds a single numerical value, -a `string_scalar` holds a single string. The data for the stored value resides in device memory. +a `string_scalar` holds a single string. -A `list_scalar` holds the underlying data of a single list. This means the underlying data can be -any type that cudf supports. For example, a `list_scalar` representing a list of integers stores a -`cudf::column` of type `INT32`. A `list_scalar` representing a list of lists of integers stores a -`cudf::column` of type `LIST`, which in turn stores a column of type `INT32`. +A scalar owns its value and validity in an Arrow-compatible, one-row `cudf::column`. This common +storage representation includes the offsets and child hierarchy required by strings and nested +types. `scalar::as_column_view()` returns an allocation-free, non-owning `scalar_column_view` +directly over that storage. |Value type|Scalar class|Notes| |-|-|-| |fixed-width|`fixed_width_scalar`| `T` can be any fixed-width type| |numeric|`numeric_scalar` | `T` can be `int8_t`, `int16_t`, `int32_t`, `int64_t`, `float` or `double`| -|fixed-point|`fixed_point_scalar` | `T` can be `numeric::decimal32` or `numeric::decimal64`| +|fixed-point|`fixed_point_scalar` | `T` can be `numeric::decimal32`, `numeric::decimal64`, or `numeric::decimal128`| |timestamp|`timestamp_scalar` | `T` can be `timestamp_D`, `timestamp_s`, etc.| |duration|`duration_scalar` | `T` can be `duration_D`, `duration_s`, etc.| |string|`string_scalar`| This class object is immutable| |list|`list_scalar`| Underlying data can be any type supported by cudf | +|struct|`struct_scalar`| Children can be any types supported by cudf | ### Construction `scalar`s can be created using either their respective constructors or using factory functions like @@ -441,15 +442,17 @@ auto s1 = static_cast(s.get()); ``` ### Passing to device -Each scalar type, except `list_scalar`, has a corresponding non-owning device view class which -allows access to the value and its validity from the device. This can be obtained using the function -`get_scalar_device_view(ScalarType s)`. Note that a device view is not provided for a base scalar -object, only for the derived typed scalar class objects. +Use `scalar::as_column_view()` and `column_device_view::create()` to access a scalar from device +code. The resulting `column_device_view` has one row, so its value and validity are accessed at +index `0`. Specialized column device views may be constructed from it for nested types. -The underlying data for `list_scalar` can be accessed via `view()` method. For non-nested data, -the device view can be obtained via function `column_device_view::create(column_view)`. For nested -data, a specialized device view for list columns can be constructed via -`lists_column_device_view(column_device_view)`. +Use `scalar::as_mutable_column_view()` and `mutable_column_device_view::create()` when device code +must modify a scalar value. Scalar validity is host-authoritative: establish it separately with +`scalar::set_valid_async()` and do not mutate the null mask through the mutable column view. + +Typed scalar device views and `get_scalar_device_view()` remain compatibility APIs. New code should +prefer column views so it works through the type-erased `scalar` interface and uses the same storage +and access patterns as columns. # libcudf Policies and Design Principles @@ -1035,8 +1038,9 @@ libcudf as well as Thrust algorithms. Most are defined in `include/detail/iterat The pair iterator is used to access elements of nullable columns as a pair containing an element's value and validity. `cudf::detail::make_pair_iterator` can be used to create a pair iterator from a -`column_device_view` or a `cudf::scalar`. `make_pair_iterator` is not available for -`mutable_column_device_view`. +`column_device_view`. Scalar callers use `scalar::as_column_view()` to create a +`column_device_view` and must retain the returned device-view owner until iterator use completes. +`make_pair_iterator` is not available for `mutable_column_device_view`. ### Null-replacement iterator diff --git a/cpp/include/cudf/ast/expressions.hpp b/cpp/include/cudf/ast/expressions.hpp index 4329f6904300..b32340e1aa67 100644 --- a/cpp/include/cudf/ast/expressions.hpp +++ b/cpp/include/cudf/ast/expressions.hpp @@ -216,10 +216,9 @@ class generic_scalar_device_view : public cudf::detail::scalar_device_view_base * * @param type The data type of the value * @param data The pointer to the data in device memory - * @param is_valid The pointer to the bool in device memory that indicates the - * validity of the stored value + * @param is_valid Pointer to the validity bitmask in device memory */ - generic_scalar_device_view(data_type type, void const* data, bool* is_valid) + generic_scalar_device_view(data_type type, void const* data, bitmask_type const* is_valid) : cudf::detail::scalar_device_view_base(type, is_valid), _data(data) { } @@ -228,11 +227,13 @@ class generic_scalar_device_view : public cudf::detail::scalar_device_view_base * * @param type The data type of the value * @param data The pointer to the data in device memory - * @param is_valid The pointer to the bool in device memory that indicates the - * validity of the stored value + * @param is_valid Pointer to the validity bitmask in device memory * @param size The size of the string in bytes */ - generic_scalar_device_view(data_type type, void const* data, bool* is_valid, size_type size) + generic_scalar_device_view(data_type type, + void const* data, + bitmask_type const* is_valid, + size_type size) : cudf::detail::scalar_device_view_base(type, is_valid), _data(data), _size(size) { } diff --git a/cpp/include/cudf/column/scalar_column_view.hpp b/cpp/include/cudf/column/scalar_column_view.hpp index 3b88f5a8a83b..d8186c9fd8d4 100644 --- a/cpp/include/cudf/column/scalar_column_view.hpp +++ b/cpp/include/cudf/column/scalar_column_view.hpp @@ -54,4 +54,45 @@ struct scalar_column_view : private column_view { [[nodiscard]] column_view const& as_column_view() const noexcept { return *this; } }; +/** + * @brief A non-owning, mutable view of one row of device column storage. + * + * This view is intended for updating scalar values. Scalar validity is managed separately by the + * owning scalar and the null mask must not be modified through this view. + * + * @ingroup column_classes + */ +struct mutable_scalar_column_view : private mutable_column_view { + /** + * @brief Construct a `mutable_scalar_column_view` from a `mutable_column_view`. + * + * @throws cudf::logic_error if the column view does not have exactly one element. + * + * @param view The mutable column view to construct from + */ + explicit mutable_scalar_column_view(mutable_column_view view) + : mutable_column_view(std::move(view)) + { + CUDF_EXPECTS( + this->size() == 1, "A scalar column view must have exactly one element.", std::logic_error); + } + + using mutable_column_view::data; + using mutable_column_view::head; + using mutable_column_view::is_empty; + using mutable_column_view::null_count; + using mutable_column_view::null_mask; + using mutable_column_view::nullable; + using mutable_column_view::offset; + using mutable_column_view::size; + using mutable_column_view::type; + + /** + * @brief Returns the underlying `mutable_column_view`. + * + * @return A reference to the underlying mutable view + */ + [[nodiscard]] mutable_column_view const& as_mutable_column_view() const noexcept { return *this; } +}; + } // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/calendrical_month_sequence.cuh b/cpp/include/cudf/detail/calendrical_month_sequence.cuh index 46004ca92efe..bc93e3960262 100644 --- a/cpp/include/cudf/detail/calendrical_month_sequence.cuh +++ b/cpp/include/cudf/detail/calendrical_month_sequence.cuh @@ -5,11 +5,11 @@ #pragma once #include +#include #include #include #include #include -#include #include #include @@ -33,8 +33,8 @@ struct calendrical_month_sequence_functor { // Return empty column if n = 0 if (n == 0) return cudf::make_empty_column(input.type()); - auto const device_input = - get_scalar_device_view(static_cast&>(const_cast(input))); + auto const input_view = input.as_column_view(); + auto const device_input = column_device_view::create(input_view.as_column_view(), stream); auto output_column_type = cudf::data_type{cudf::type_to_id()}; auto output = cudf::make_fixed_width_column( output_column_type, n, cudf::mask_state::UNALLOCATED, stream, mr); @@ -43,9 +43,9 @@ struct calendrical_month_sequence_functor { cuda::counting_iterator{0}, cuda::counting_iterator{n}, output->mutable_view().begin(), - [initial = device_input, months] __device__(size_type i) { + [initial = *device_input, months] __device__(size_type i) { return datetime::detail::add_calendrical_months_with_scale_back( - initial.value(), cuda::std::chrono::months{i * months}); + initial.element(0), cuda::std::chrono::months{i * months}); }); return output; diff --git a/cpp/include/cudf/detail/iterator.cuh b/cpp/include/cudf/detail/iterator.cuh index 797bc0acada8..7141e1e6efd0 100644 --- a/cpp/include/cudf/detail/iterator.cuh +++ b/cpp/include/cudf/detail/iterator.cuh @@ -23,8 +23,6 @@ #pragma once #include -#include -#include #include #include @@ -316,335 +314,5 @@ CUDF_HOST_DEVICE auto inline make_validity_iterator(column_device_view const& co return make_counting_transform_iterator(cudf::size_type{0}, validity_accessor{column}); } -/** - * @brief Constructs a constant device iterator over a scalar's validity. - * - * Dereferencing the returned iterator returns a `bool`. - * - * For `p = *(iter + i)`, `p` is the validity of the scalar. - * - * @tparam safe unused. This template parameter exists to enforce the same - * template interface as @ref make_validity_iterator(column_device_view const&). - * @param scalar_value The scalar to iterate - * @return auto Iterator that returns scalar validity - */ -template -auto inline make_validity_iterator(scalar const& scalar_value) -{ - return cuda::make_constant_iterator(scalar_value.is_valid()); -} - -/** - * @brief value accessor for scalar with valid data. - * The unary functor returns data of Element type of the scalar. - * - * @throws `cudf::logic_error` if scalar datatype and Element type mismatch. - * - * @tparam Element The type of return type of functor - */ -template -struct scalar_value_accessor { - using ScalarType = scalar_type_t; - using ScalarDeviceType = scalar_device_type_t; - ScalarDeviceType const dscalar; ///< scalar device view - - scalar_value_accessor(scalar const& scalar_value) - : dscalar(get_scalar_device_view(static_cast(const_cast(scalar_value)))) - { - CUDF_EXPECTS(type_id_matches_device_storage_type(scalar_value.type().id()), - "the data type mismatch"); - } - - __device__ inline Element const operator()(size_type) const { return dscalar.value(); } -}; - -/** - * @brief Constructs a constant device iterator over a scalar's value. - * - * Dereferencing the returned iterator returns a `Element`. - * - * For `p = *(iter + i)`, `p` is the value stored in the scalar. - * - * The behavior is undefined if the scalar is destroyed before iterator dereferencing. - * - * @throws cudf::logic_error if scalar datatype and Element type mismatch. - * @throws cudf::logic_error if scalar is null. - * @throws cudf::logic_error if the returned iterator is dereferenced in host - * - * @tparam Element The type of element in the scalar - * @param scalar_value The scalar to iterate - * @return auto Iterator that returns scalar value - */ -template -auto inline make_scalar_iterator(scalar const& scalar_value) -{ - CUDF_EXPECTS(data_type(type_to_id()) == scalar_value.type(), "the data type mismatch"); - CUDF_EXPECTS(scalar_value.is_valid(), "the scalar value must be valid"); - return cuda::transform_iterator(cuda::make_constant_iterator(0), - scalar_value_accessor{scalar_value}); -} - -/** - * @brief Optional accessor for a scalar - * - * The `scalar_optional_accessor` always returns a `cuda::std::optional` of the scalar. - * The validity of the optional is determined by the `Nullate` parameter which may - * be one of the following: - * - * - `nullate::YES` means that the scalar may be valid or invalid and the optional returned - * will contain a value only if the scalar is valid. - * - * - `nullate::NO` means the caller attests that the scalar will always be valid, - * no checks will occur and `cuda::std::optional{column[i]}` will return a value - * for each `i`. - * - * - `nullate::DYNAMIC` defers the assumption of nullability to runtime and the caller - * specifies if the scalar may be valid or invalid. - * For `DYNAMIC{true}` the return value will be a `cuda::std::optional{scalar}` when the - * scalar is valid and a `cuda::std::optional{}` when the scalar is invalid. - * For `DYNAMIC{false}` the return value will always be a `cuda::std::optional{scalar}`. - * - * @throws `cudf::logic_error` if scalar datatype and Element type mismatch. - * - * @tparam Element The type of return type of functor - * @tparam Nullate A cudf::nullate type describing how to check for nulls. - */ -template -struct scalar_optional_accessor : public scalar_value_accessor { - using super_t = scalar_value_accessor; - using value_type = cuda::std::optional; - - scalar_optional_accessor(scalar const& scalar_value, Nullate with_nulls) - : scalar_value_accessor(scalar_value), has_nulls{with_nulls} - { - } - - __device__ inline value_type const operator()(size_type) const - { - if (has_nulls && !super_t::dscalar.is_valid()) { return value_type{cuda::std::nullopt}; } - - if constexpr (cudf::is_fixed_point()) { - using namespace numeric; - using rep = typename Element::rep; - auto const value = super_t::dscalar.rep(); - auto const scale = scale_type{super_t::dscalar.type().scale()}; - return Element{scaled_integer{value, scale}}; - } else { - return Element{super_t::dscalar.value()}; - } - } - - Nullate has_nulls{}; -}; - -/** - * @brief pair accessor for scalar. - * The unary functor returns a pair of data of Element type and bool validity of the scalar. - * - * @throws `cudf::logic_error` if scalar datatype and Element type mismatch. - * - * @tparam Element The type of return type of functor - */ -template -struct scalar_pair_accessor : public scalar_value_accessor { - using super_t = scalar_value_accessor; - using value_type = cuda::std::pair; - scalar_pair_accessor(scalar const& scalar_value) : scalar_value_accessor(scalar_value) {} - - __device__ inline value_type const operator()(size_type) const - { - return {Element(super_t::dscalar.value()), super_t::dscalar.is_valid()}; - } -}; - -/** - * @brief Utility to discard template type arguments. - * - * Substitute for std::void_t. - * - * @tparam T Ignored template parameter - */ -template -using void_t = void; - -/** - * @brief Compile-time reflection to check if `Element` type has a `rep()` member. - */ -template -struct has_rep_member : std::false_type {}; - -template -struct has_rep_member().rep())>> : std::true_type {}; - -/** - * @brief Pair accessor for scalar's representation value and validity. - * - * @tparam Element The type of element in the scalar. - */ -template -struct scalar_representation_pair_accessor : public scalar_value_accessor { - using base = scalar_value_accessor; - using rep_type = device_storage_type_t; - using value_type = cuda::std::pair; - - scalar_representation_pair_accessor(scalar const& scalar_value) : base(scalar_value) {} - - __device__ inline value_type const operator()(size_type) const - { - return {get_rep(base::dscalar), base::dscalar.is_valid()}; - } - - private: - template - __device__ inline rep_type get_rep(DeviceScalar const& dscalar) const - requires(!has_rep_member::value) - { - return dscalar.value(); - } - - template - __device__ inline rep_type get_rep(DeviceScalar const& dscalar) const - requires(has_rep_member::value) - { - return dscalar.rep(); - } -}; - -/** - * @brief Constructs an optional iterator over a scalar's values and its validity. - * - * Dereferencing the returned iterator returns a `cuda::std::optional`. - * - * The element of this iterator contextually converts to bool. The conversion returns true - * if the object contains a value and false if it does not contain a value. - * - * The iterator behavior is undefined if the scalar is destroyed before iterator dereferencing. - * - * Calling this function with `nullate::DYNAMIC` defers the assumption - * of nullability to runtime with the caller indicating if the scalar is valid. - * - * @code{.cpp} - * template - * void some_function(cudf::column_view const& col_view, - * scalar const& scalar_value, - * bool col_has_nulls){ - * auto d_col = cudf::column_device_view::create(col_view); - * auto column_iterator = cudf::detail::make_optional_iterator( - * d_col, cudf::nullate::DYNAMIC{col_has_nulls}); - * auto scalar_iterator = cudf::detail::make_optional_iterator( - * scalar_value, cudf::nullate::DYNAMIC{scalar_value.is_valid()}); - * //use iterators - * } - * @endcode - * - * Calling this function with `nullate::YES` means that the scalar maybe invalid - * and the optional return might not contain a value. - * Calling this function with `nullate::NO` means that the scalar is valid - * and the optional returned will always contain a value. - * - * @code{.cpp} - * template - * void some_function(cudf::column_view const& col_view, scalar const& scalar_value){ - * auto d_col = cudf::column_device_view::create(col_view); - * if constexpr(any_nulls) { - * auto column_iterator = - * cudf::detail::make_optional_iterator(d_col, cudf::nullate::YES{}); - * auto scalar_iterator = - * cudf::detail::make_optional_iterator(scalar_value, cudf::nullate::YES{}); - * //use iterators - * } else { - * auto column_iterator = - * cudf::detail::make_optional_iterator(d_col, cudf::nullate::NO{}); - * auto scalar_iterator = - * cudf::detail::make_optional_iterator(scalar_value, cudf::nullate::NO{}); - * //use iterators - * } - * } - * @endcode - * - * @throws cudf::logic_error if scalar datatype and Element type mismatch. - * - * @tparam Element The type of elements in the scalar - * @tparam Nullate A cudf::nullate type describing how to check for nulls. - * - * @param scalar_value The scalar to be returned by the iterator. - * @param has_nulls Indicates if the scalar value may be invalid. - * @return Iterator that returns scalar and the validity of the scalar in a cuda::std::optional - */ -template -auto inline make_optional_iterator(scalar const& scalar_value, Nullate has_nulls) -{ - CUDF_EXPECTS(type_id_matches_device_storage_type(scalar_value.type().id()), - "the data type mismatch"); - return cuda::transform_iterator( - cuda::make_constant_iterator(0), - scalar_optional_accessor{scalar_value, has_nulls}); -} - -/** - * @brief Constructs a constant device pair iterator over a scalar's value and its validity. - * - * Dereferencing the returned iterator returns a `cuda::std::pair`. - * - * If scalar is valid, then for `p = *(iter + i)`, `p.first` contains - * the value of the scalar and `p.second == true`. - * - * Else, if the scalar is null, then the value of `p.first` is undefined and `p.second == false`. - * - * The behavior is undefined if the scalar is destroyed before iterator dereferencing. - * - * @throws cudf::logic_error if scalar datatype and Element type mismatch. - * @throws cudf::logic_error if the returned iterator is dereferenced in host - * - * @tparam Element The type of elements in the scalar - * @tparam bool unused. This template parameter exists to enforce same - * template interface as @ref make_pair_iterator(column_device_view const&). - * @param scalar_value The scalar to iterate - * @return auto Iterator that returns scalar, and validity of the scalar in a pair - */ -template -auto inline make_pair_iterator(scalar const& scalar_value) -{ - CUDF_EXPECTS(type_id_matches_device_storage_type(scalar_value.type().id()), - "the data type mismatch"); - return cuda::transform_iterator(cuda::make_constant_iterator(0), - scalar_pair_accessor{scalar_value}); -} - -/** - * @brief Constructs a constant device pair iterator over a scalar's representative value - * and its validity. - * - * Dereferencing the returned iterator returns a `cuda::std::pair`. - * E.g. For a valid `decimal32` row, a `cuda::std::pair` is returned, - * with the value set to the `int32_t` representative value of the decimal, - * and validity `true`, indicating that the row is valid. - * - * If scalar is valid, then for `p = *(iter + i)`, `p.first` contains - * the representative value of the scalar and `p.second == true`. - * - * Else, if the scalar is null, then the value of `p.first` is undefined and `p.second == false`. - * - * The behavior is undefined if the scalar is destroyed before iterator dereferencing. - * - * @throws cudf::logic_error if scalar datatype and Element type mismatch. - * @throws cudf::logic_error if the returned iterator is dereferenced in host - * - * @tparam Element The type of elements in the scalar - * @tparam bool unused. This template parameter exists to enforce same - * template interface as @ref make_pair_iterator(column_device_view const&). - * @param scalar_value The scalar to iterate - * @return auto Iterator that returns scalar's representative value, - * and validity of the scalar in a pair - */ -template -auto make_pair_rep_iterator(scalar const& scalar_value) -{ - CUDF_EXPECTS(type_id_matches_device_storage_type(scalar_value.type().id()), - "the data type mismatch"); - return make_counting_transform_iterator( - 0, scalar_representation_pair_accessor{scalar_value}); -} - } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/scalar/scalar.hpp b/cpp/include/cudf/scalar/scalar.hpp index c822c6e934de..5d1b61d05b44 100644 --- a/cpp/include/cudf/scalar/scalar.hpp +++ b/cpp/include/cudf/scalar/scalar.hpp @@ -5,7 +5,9 @@ #pragma once #include +#include #include +#include #include #include #include @@ -17,6 +19,7 @@ #include +#include #include #include @@ -56,40 +59,53 @@ class scalar { /** * @brief Updates the validity of the value. * + * Updates the host null-count metadata immediately and enqueues the corresponding device + * bitmask update on `stream`. + * * @param is_valid true: set the value to valid. false: set it to null. - * @param stream CUDA stream used for device memory operations. + * @param stream CUDA stream used for the device bitmask update. */ void set_valid_async(bool is_valid, cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Indicates whether the scalar contains a valid value. * - * @note Using the value when `is_valid() == false` is undefined behavior. In addition, this - * function does a stream synchronization. + * This reads host null-count metadata and does not inspect or synchronize the device bitmask. * - * @param stream CUDA stream used for device memory operations. + * @note Using the value when `is_valid() == false` is undefined behavior. + * + * @param stream Retained for API compatibility; no synchronization is performed. * @return true Value is valid * @return false Value is invalid/null */ [[nodiscard]] bool is_valid(cuda::stream_ref stream = cudf::get_default_stream()) const; /** - * @brief Returns a raw pointer to the validity bool in device memory. + * @brief Return a const raw pointer to the validity bitmask in device memory. * - * @return Raw pointer to the validity bool in device memory + * @return Raw pointer to the validity bitmask in device memory */ - bool* validity_data(); + [[nodiscard]] bitmask_type const* validity_data() const; /** - * @brief Return a const raw pointer to the validity bool in device memory. + * @brief Returns an allocation-free one-row column view of this scalar. * - * @return Raw pointer to the validity bool in device memory + * @return A view directly referencing the scalar's owned column storage */ - [[nodiscard]] bool const* validity_data() const; + [[nodiscard]] scalar_column_view as_column_view() const; + + /** + * @brief Returns an allocation-free mutable one-row column view of this scalar. + * + * Device code may modify the value through this view. Validity must be established separately + * with `set_valid_async()`; the null mask must not be modified through this view. + * + * @return A mutable view directly referencing the scalar's owned column storage + */ + [[nodiscard]] mutable_scalar_column_view as_mutable_column_view(); protected: - data_type _type{type_id::EMPTY}; ///< Logical type of value in the scalar - cudf::detail::device_scalar _is_valid; ///< Device bool signifying validity + cudf::column _storage; ///< Arrow-compatible one-row column storage /** * @brief Move constructor for scalar. @@ -123,6 +139,13 @@ class scalar { bool is_valid = false, cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + + /** + * @brief Construct a scalar by taking ownership of a one-row column. + * + * @param storage One-row column storage + */ + explicit scalar(cudf::column&& storage); }; namespace detail { @@ -190,7 +213,7 @@ class fixed_width_scalar : public scalar { [[nodiscard]] T const* data() const; protected: - cudf::detail::device_scalar _data; ///< device memory containing the value + mutable cudf::detail::host_vector _bounce_buffer; ///< Host staging for async value updates /** * @brief Construct a new fixed width scalar object. @@ -400,9 +423,6 @@ class fixed_point_scalar : public scalar { * @return a const raw pointer to the value in device memory */ [[nodiscard]] rep_type const* data() const; - - protected: - cudf::detail::device_scalar _data; ///< device memory containing the value }; /** @@ -525,8 +545,24 @@ class string_scalar : public scalar { */ [[nodiscard]] char const* data() const; - protected: - rmm::device_buffer _data{}; ///< device memory containing the string + private: + struct string_storage; + + explicit string_scalar(string_storage&& storage); + + static string_storage make_storage(std::string_view string, + bool is_valid, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr); + + string_scalar(rmm::device_buffer&& data, + size_type size, + bool is_valid, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr); + + size_type _size{}; ///< Host metadata for the number of string bytes + std::optional> _host_data; ///< Async host-source staging }; /** @@ -760,9 +796,6 @@ class list_scalar : public scalar { * @return A non-owning, immutable view to underlying device data */ [[nodiscard]] column_view view() const; - - private: - cudf::column _data; }; /** @@ -845,8 +878,6 @@ class struct_scalar : public scalar { [[nodiscard]] table_view view() const; private: - table _data; - /** * @brief Check if all the input columns constructing this struct scalar have valid size. */ diff --git a/cpp/include/cudf/scalar/scalar_device_view.cuh b/cpp/include/cudf/scalar/scalar_device_view.cuh index ce722312f227..8538ed6b530a 100644 --- a/cpp/include/cudf/scalar/scalar_device_view.cuh +++ b/cpp/include/cudf/scalar/scalar_device_view.cuh @@ -7,6 +7,7 @@ #include #include #include +#include /** * @file scalar_device_view.cuh @@ -36,29 +37,23 @@ class scalar_device_view_base { * @return true The element is valid * @return false The element is null */ - [[nodiscard]] __device__ bool is_valid() const noexcept { return *_is_valid; } - - /** - * @brief Updates the validity of the value - * - * @param is_valid true: set the value to valid. false: set it to null - */ - __device__ void set_valid(bool is_valid) noexcept { *_is_valid = is_valid; } + [[nodiscard]] __device__ bool is_valid() const noexcept { return bit_is_set(_is_valid, 0); } protected: data_type _type{type_id::EMPTY}; ///< Value data type - bool* _is_valid{}; ///< Pointer to device memory containing - ///< boolean representing validity of the value. + bitmask_type const* _is_valid{}; ///< Pointer to the scalar validity bitmask. /** - * @brief Construct a new scalar device view base object from a device pointer - * and a validity boolean. + * @brief Construct a new scalar device view base object from a device pointer + * and a validity bitmask. * * @param type The data type of the scalar - * @param is_valid Pointer to device memory containing boolean representing - * validity of the scalar. + * @param is_valid Pointer to the validity bitmask in device memory. */ - scalar_device_view_base(data_type type, bool* is_valid) : _type(type), _is_valid(is_valid) {} + scalar_device_view_base(data_type type, bitmask_type const* is_valid) + : _type(type), _is_valid(is_valid) + { + } scalar_device_view_base() = default; }; @@ -139,10 +134,9 @@ class fixed_width_scalar_device_view_base : public detail::scalar_device_view_ba * * @param type The data type of the value * @param data The pointer to the data in device memory - * @param is_valid The pointer to the bool in device memory that indicates the - * validity of the stored value + * @param is_valid Pointer to the validity bitmask in device memory */ - fixed_width_scalar_device_view_base(data_type type, void* data, bool* is_valid) + fixed_width_scalar_device_view_base(data_type type, void* data, bitmask_type const* is_valid) : detail::scalar_device_view_base(type, is_valid), _data(data) { } @@ -206,10 +200,9 @@ class fixed_width_scalar_device_view : public detail::fixed_width_scalar_device_ * * @param type The data type of the value * @param data The pointer to the data in device memory - * @param is_valid The pointer to the bool in device memory that indicates the - * validity of the stored value + * @param is_valid Pointer to the validity bitmask in device memory */ - fixed_width_scalar_device_view(data_type type, T* data, bool* is_valid) + fixed_width_scalar_device_view(data_type type, T* data, bitmask_type const* is_valid) : detail::fixed_width_scalar_device_view_base(type, data, is_valid) { } @@ -228,10 +221,9 @@ class numeric_scalar_device_view : public detail::fixed_width_scalar_device_view * * @param type The data type of the value * @param data The pointer to the data in device memory - * @param is_valid The pointer to the bool in device memory that indicates the - * validity of the stored value + * @param is_valid Pointer to the validity bitmask in device memory */ - numeric_scalar_device_view(data_type type, T* data, bool* is_valid) + numeric_scalar_device_view(data_type type, T* data, bitmask_type const* is_valid) : detail::fixed_width_scalar_device_view(type, data, is_valid) { } @@ -250,10 +242,9 @@ class fixed_point_scalar_device_view : public detail::scalar_device_view_base { * * @param type The data type of the value * @param data The pointer to the data in device memory - * @param is_valid The pointer to the bool in device memory that indicates the - * validity of the stored value + * @param is_valid Pointer to the validity bitmask in device memory */ - fixed_point_scalar_device_view(data_type type, rep_type* data, bool* is_valid) + fixed_point_scalar_device_view(data_type type, rep_type* data, bitmask_type const* is_valid) : detail::scalar_device_view_base(type, is_valid), _data(data) { } @@ -289,11 +280,13 @@ class string_scalar_device_view : public detail::scalar_device_view_base { * * @param type The data type of the value * @param data The pointer to the string data in device memory - * @param is_valid The pointer to the bool in device memory that indicates the - * validity of the stored value + * @param is_valid Pointer to the validity bitmask in device memory * @param size The pointer to the size of the string in device memory */ - string_scalar_device_view(data_type type, char const* data, bool* is_valid, size_type size) + string_scalar_device_view(data_type type, + char const* data, + bitmask_type const* is_valid, + size_type size) : detail::scalar_device_view_base(type, is_valid), _data(data), _size(size) { } @@ -341,10 +334,9 @@ class timestamp_scalar_device_view : public detail::fixed_width_scalar_device_vi * * @param type The data type of the value * @param data The pointer to the data in device memory - * @param is_valid The pointer to the bool in device memory that indicates the - * validity of the stored value + * @param is_valid Pointer to the validity bitmask in device memory */ - timestamp_scalar_device_view(data_type type, T* data, bool* is_valid) + timestamp_scalar_device_view(data_type type, T* data, bitmask_type const* is_valid) : detail::fixed_width_scalar_device_view(type, data, is_valid) { } @@ -361,10 +353,9 @@ class duration_scalar_device_view : public detail::fixed_width_scalar_device_vie * * @param type The data type of the value * @param data The pointer to the data in device memory - * @param is_valid The pointer to the bool in device memory that indicates the - * validity of the stored value + * @param is_valid Pointer to the validity bitmask in device memory */ - duration_scalar_device_view(data_type type, T* data, bool* is_valid) + duration_scalar_device_view(data_type type, T* data, bitmask_type const* is_valid) : detail::fixed_width_scalar_device_view(type, data, is_valid) { } diff --git a/cpp/include/cudf/scalar/scalar_factories.hpp b/cpp/include/cudf/scalar/scalar_factories.hpp index efc1de56b64d..81917ee047ac 100644 --- a/cpp/include/cudf/scalar/scalar_factories.hpp +++ b/cpp/include/cudf/scalar/scalar_factories.hpp @@ -32,7 +32,7 @@ namespace CUDF_EXPORT cudf { * * @param type The desired numeric element type * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @returns An uninitialized numeric scalar */ std::unique_ptr make_numeric_scalar( @@ -49,7 +49,7 @@ std::unique_ptr make_numeric_scalar( * * @param type The desired timestamp element type * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @return An uninitialized timestamp scalar */ std::unique_ptr make_timestamp_scalar( @@ -66,7 +66,7 @@ std::unique_ptr make_timestamp_scalar( * * @param type The desired duration element type * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @return An uninitialized duration scalar */ std::unique_ptr make_duration_scalar( @@ -83,7 +83,7 @@ std::unique_ptr make_duration_scalar( * * @param type The desired fixed-width element type * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @return An uninitialized fixed-width scalar */ std::unique_ptr make_fixed_width_scalar( @@ -100,7 +100,7 @@ std::unique_ptr make_fixed_width_scalar( * * @param string The `std::string` to copy to device * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @returns A string scalar with the contents of `string` */ std::unique_ptr make_string_scalar( @@ -115,7 +115,7 @@ std::unique_ptr make_string_scalar( * * @param type The desired element type * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @returns A scalar of type `type` */ std::unique_ptr make_default_constructed_scalar( @@ -130,7 +130,7 @@ std::unique_ptr make_default_constructed_scalar( * * @param input Immutable view of input column to emulate * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @returns A scalar of type of `input` column */ std::unique_ptr make_empty_scalar_like( @@ -144,7 +144,7 @@ std::unique_ptr make_empty_scalar_like( * @tparam T Datatype of the value to be represented by the scalar * @param value The value to store in the scalar object * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @returns A scalar of type `T` */ template @@ -163,7 +163,7 @@ std::unique_ptr make_fixed_width_scalar( * @param value The value to store in the scalar object * @param scale The scale of the fixed point value * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @returns A scalar of type `T` */ template @@ -181,7 +181,7 @@ std::unique_ptr make_fixed_point_scalar( * * @param elements Elements of the list * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @returns A list scalar */ std::unique_ptr make_list_scalar( @@ -196,7 +196,7 @@ std::unique_ptr make_list_scalar( * * @param data The columnar data to store in the scalar object * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @returns A struct scalar */ std::unique_ptr make_struct_scalar( @@ -211,7 +211,7 @@ std::unique_ptr make_struct_scalar( * * @param data The columnar data to store in the scalar object * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + * @param mr Device memory resource used to allocate the scalar's column storage. * @returns A struct scalar */ std::unique_ptr make_struct_scalar( diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu index 0e639d90e95d..c6bb46fbf100 100644 --- a/cpp/src/binaryop/compiled/binary_ops.cu +++ b/cpp/src/binaryop/compiled/binary_ops.cu @@ -10,10 +10,7 @@ #include #include #include -#include #include -#include -#include #include #include @@ -29,183 +26,60 @@ namespace binops { namespace compiled { namespace { -/** - * @brief Converts scalar to column_view with single element. - * - * @return pair with column_view and column containing any auxiliary data to create column_view from - * scalar - */ -struct scalar_as_column_view { - using return_type = typename std::pair>; - template ())> - return_type operator()(scalar const& s, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) - { - auto& h_scalar_type_view = static_cast&>(const_cast(s)); - - // Valid scalar needs no null mask - if (s.is_valid(stream)) { - auto col_v = column_view(s.type(), 1, h_scalar_type_view.data(), nullptr, 0); - return std::pair{col_v, nullptr}; - } - - // Null scalar needs a single-element null mask (kept alive by an auxiliary column) as its - // validity (bool) cannot be reinterpreted as `bitmask_type` without reading out of bounds. - auto null_mask = cudf::detail::create_null_mask(1, cudf::mask_state::ALL_NULL, stream, mr); - auto const* null_mask_ptr = static_cast(null_mask.data()); - auto aux_col = std::make_unique( - data_type{type_id::INT8}, 0, rmm::device_buffer{}, std::move(null_mask), 1); - auto col_v = column_view(s.type(), 1, h_scalar_type_view.data(), null_mask_ptr, 1); - return std::pair{col_v, std::move(aux_col)}; - } - template ())> - return_type operator()(scalar const&, cuda::stream_ref, rmm::device_async_resource_ref) - { - CUDF_FAIL("Unsupported type"); - } -}; -// specialization for cudf::string_view -template <> -scalar_as_column_view::return_type scalar_as_column_view::operator()( - scalar const& s, cuda::stream_ref stream, rmm::device_async_resource_ref mr) -{ - using T = cudf::string_view; - auto& h_scalar_type_view = static_cast&>(const_cast(s)); - - // build offsets column from the string size - auto offsets_transformer_itr = cuda::make_constant_iterator(h_scalar_type_view.size()); - auto offsets_column = std::get<0>(cudf::detail::make_offsets_child_column( - offsets_transformer_itr, offsets_transformer_itr + 1, stream, mr)); - - // Valid scalar needs no null mask. The offsets child column is kept alive to back the returned - // string column_view. - if (s.is_valid(stream)) { - auto col_v = cudf::column_view( - s.type(), 1, h_scalar_type_view.data(), nullptr, 0, 0, {offsets_column->view()}); - return std::pair{col_v, std::move(offsets_column)}; - } - - // Null scalar needs a single-element null mask and offsets (kepy alive by an auxiliary column) as - // its validity (bool) cannot be reinterpreted as `bitmask_type` without reading out of bounds. - auto null_mask = cudf::detail::create_null_mask(1, cudf::mask_state::ALL_NULL, stream, mr); - auto const* null_mask_ptr = static_cast(null_mask.data()); - auto col_v = cudf::column_view( - s.type(), 1, h_scalar_type_view.data(), null_mask_ptr, 1, 0, {offsets_column->view()}); - std::vector> children; - children.push_back(std::move(offsets_column)); - auto aux_col = std::make_unique(data_type{type_id::INT8}, - 0, - rmm::device_buffer{}, - std::move(null_mask), - 1, - std::move(children)); - return std::pair{col_v, std::move(aux_col)}; -} - -// specializing for struct column -template <> -scalar_as_column_view::return_type scalar_as_column_view::operator()( - scalar const& s, cuda::stream_ref stream, rmm::device_async_resource_ref mr) -{ - auto col = make_column_from_scalar(s, 1, stream, mr); - return std::pair{col->view(), std::move(col)}; -} - -/** - * @brief Converts scalar to column_view with single element. - * - * @param scal scalar to convert - * @param stream CUDA stream used for device memory operations and kernel launches. - * @param mr Device memory resource used to allocate the returned column's device memory - * @return pair with column_view and column containing any auxiliary data to create - * column_view from scalar - */ -auto scalar_to_column_view( - scalar const& scal, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) -{ - return type_dispatcher(scal.type(), scalar_as_column_view{}, scal, stream, mr); -} - -// This functor does the actual comparison between string column value and a scalar string -// or between two string column values using a comparator -template +// This functor compares string inputs represented by column device views. Scalar inputs use row 0. +template struct compare_functor { - LhsDeviceViewT const lhs_dev_view_; // Scalar or a column device view - lhs - RhsDeviceViewT const rhs_dev_view_; // Scalar or a column device view - rhs + column_device_view const lhs_dev_view_; + column_device_view const rhs_dev_view_; CompareFunc const cfunc_; // Comparison function - - compare_functor(LhsDeviceViewT const& lhs_dev_view, - RhsDeviceViewT const& rhs_dev_view, - CompareFunc cf) - : lhs_dev_view_(lhs_dev_view), rhs_dev_view_(rhs_dev_view), cfunc_(cf) - { - } - - // This is used to compare a scalar and a column value - template - __device__ inline OutT operator()(cudf::size_type i) const - requires(std::is_same_v && - !std::is_same_v) - { - return cfunc_(lhs_dev_view_.is_valid(i), - rhs_dev_view_.is_valid(), - lhs_dev_view_.is_valid(i) ? lhs_dev_view_.template element(i) - : cudf::string_view{}, - rhs_dev_view_.is_valid() ? rhs_dev_view_.value() : cudf::string_view{}); - } - - // This is used to compare a scalar and a column value - template - __device__ inline OutT operator()(cudf::size_type i) const - requires(!std::is_same_v && - std::is_same_v) + bool const is_lhs_scalar_; + bool const is_rhs_scalar_; + + compare_functor(column_device_view const& lhs_dev_view, + column_device_view const& rhs_dev_view, + CompareFunc cf, + bool is_lhs_scalar, + bool is_rhs_scalar) + : lhs_dev_view_(lhs_dev_view), + rhs_dev_view_(rhs_dev_view), + cfunc_(cf), + is_lhs_scalar_(is_lhs_scalar), + is_rhs_scalar_(is_rhs_scalar) { - return cfunc_(lhs_dev_view_.is_valid(), - rhs_dev_view_.is_valid(i), - lhs_dev_view_.is_valid() ? lhs_dev_view_.value() : cudf::string_view{}, - rhs_dev_view_.is_valid(i) ? rhs_dev_view_.template element(i) - : cudf::string_view{}); } - // This is used to compare 2 column values - template __device__ inline OutT operator()(cudf::size_type i) const - requires(std::is_same_v && - std::is_same_v) { - return cfunc_(lhs_dev_view_.is_valid(i), - rhs_dev_view_.is_valid(i), - lhs_dev_view_.is_valid(i) ? lhs_dev_view_.template element(i) - : cudf::string_view{}, - rhs_dev_view_.is_valid(i) ? rhs_dev_view_.template element(i) - : cudf::string_view{}); + auto const lhs_index = is_lhs_scalar_ ? 0 : i; + auto const rhs_index = is_rhs_scalar_ ? 0 : i; + return cfunc_( + lhs_dev_view_.is_valid(lhs_index), + rhs_dev_view_.is_valid(rhs_index), + lhs_dev_view_.is_valid(lhs_index) + ? lhs_dev_view_.template element(lhs_index) + : cudf::string_view{}, + rhs_dev_view_.is_valid(rhs_index) + ? rhs_dev_view_.template element(rhs_index) + : cudf::string_view{}); } }; // This functor performs null aware binop between two columns or a column and a scalar by // iterating over them on the device struct null_considering_binop { - [[nodiscard]] auto get_device_view(cudf::scalar const& scalar_item) const - { - return get_scalar_device_view( - static_cast&>(const_cast(scalar_item))); - } - - [[nodiscard]] auto get_device_view(column_device_view const& col_item) const { return col_item; } - - template - void populate_out_col(LhsViewT const& lhsv, - RhsViewT const& rhsv, + template + void populate_out_col(column_device_view const& lhsv, + column_device_view const& rhsv, cudf::size_type col_size, cuda::stream_ref stream, CompareFunc cfunc, + bool is_lhs_scalar, + bool is_rhs_scalar, OutT* out_col) const { // Create binop functor instance - compare_functor binop_func{lhsv, rhsv, cfunc}; + compare_functor binop_func{ + lhsv, rhsv, cfunc, is_lhs_scalar, is_rhs_scalar}; // Execute it on every element thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), @@ -216,18 +90,16 @@ struct null_considering_binop { } // This is invoked to perform comparison between cudf string types - template - std::unique_ptr operator()(LhsT const& lhs, - RhsT const& rhs, + std::unique_ptr operator()(column_device_view const& lhs, + column_device_view const& rhs, binary_operator op, data_type output_type, cudf::size_type col_size, + bool is_lhs_scalar, + bool is_rhs_scalar, cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { - // Create device views for inputs - auto const lhs_dev_view = get_device_view(lhs); - auto const rhs_dev_view = get_device_view(rhs); // Validate input CUDF_EXPECTS(output_type.id() == lhs.type().id(), "Output column type should match input column type"); @@ -255,8 +127,14 @@ 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()); + populate_out_col(lhs, + rhs, + col_size, + stream, + minmax_func, + is_lhs_scalar, + is_rhs_scalar, + out_col_strings.data()); // Create an output column with the resultant strings return cudf::make_strings_column(out_col_strings, invalid_str, stream, mr); @@ -278,8 +156,11 @@ std::unique_ptr string_null_min_max(scalar const& lhs, CUDF_EXPECTS(op == binary_operator::NULL_MAX or op == binary_operator::NULL_MIN, "Unsupported binary operation"); if (rhs.is_empty()) return cudf::make_empty_column(output_type); + auto lhs_view = lhs.as_column_view(); + auto lhs_device_view = cudf::column_device_view::create(lhs_view.as_column_view(), stream); auto rhs_device_view = cudf::column_device_view::create(rhs, stream); - return null_considering_binop{}(lhs, *rhs_device_view, op, output_type, rhs.size(), stream, mr); + return null_considering_binop{}( + *lhs_device_view, *rhs_device_view, op, output_type, rhs.size(), true, false, stream, mr); } std::unique_ptr string_null_min_max(column_view const& lhs, @@ -296,7 +177,10 @@ std::unique_ptr string_null_min_max(column_view const& lhs, "Unsupported binary operation"); if (lhs.is_empty()) return cudf::make_empty_column(output_type); auto lhs_device_view = cudf::column_device_view::create(lhs, stream); - return null_considering_binop{}(*lhs_device_view, rhs, op, output_type, lhs.size(), stream, mr); + auto rhs_view = rhs.as_column_view(); + auto rhs_device_view = cudf::column_device_view::create(rhs_view.as_column_view(), stream); + return null_considering_binop{}( + *lhs_device_view, *rhs_device_view, op, output_type, lhs.size(), false, true, stream, mr); } std::unique_ptr string_null_min_max(column_view const& lhs, @@ -316,7 +200,7 @@ std::unique_ptr string_null_min_max(column_view const& lhs, auto lhs_device_view = cudf::column_device_view::create(lhs, stream); auto rhs_device_view = cudf::column_device_view::create(rhs, stream); return null_considering_binop{}( - *lhs_device_view, *rhs_device_view, op, output_type, lhs.size(), stream, mr); + *lhs_device_view, *rhs_device_view, op, output_type, lhs.size(), false, false, stream, mr); } void operator_dispatcher(mutable_column_view& out, @@ -388,8 +272,8 @@ void binary_operation(mutable_column_view& out, binary_operator op, cuda::stream_ref stream) { - auto [lhsv, aux] = scalar_to_column_view(lhs, stream); - operator_dispatcher(out, lhsv, rhs, true, false, op, stream); + auto const lhs_view = lhs.as_column_view(); + operator_dispatcher(out, lhs_view.as_column_view(), rhs, true, false, op, stream); } // vector_scalar void binary_operation(mutable_column_view& out, @@ -398,8 +282,8 @@ void binary_operation(mutable_column_view& out, binary_operator op, cuda::stream_ref stream) { - auto [rhsv, aux] = scalar_to_column_view(rhs, stream); - operator_dispatcher(out, lhs, rhsv, false, true, op, stream); + auto const rhs_view = rhs.as_column_view(); + operator_dispatcher(out, lhs, rhs_view.as_column_view(), false, true, op, stream); } namespace detail { diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu index 391b6c59130a..a31e97b4fc11 100644 --- a/cpp/src/copying/copy.cu +++ b/cpp/src/copying/copy.cu @@ -52,13 +52,25 @@ struct get_iterable_device_view { } template - auto operator()(T const& input, cuda::stream_ref) + auto operator()(T const& input, cuda::stream_ref stream) requires(std::is_same_v) { - return &input; + auto const input_view = input.as_column_view(); + return cudf::column_device_view::create(input_view.as_column_view(), stream); } }; +template +auto make_iterable_optional_iterator(column_device_view const& input, bool nullable) +{ + auto values = cudf::detail::make_optional_iterator(input, nullate::DYNAMIC{nullable}); + if constexpr (std::is_same_v) { + return cuda::make_permutation_iterator(values, cuda::make_constant_iterator(0)); + } else { + return values; + } +} + template struct copy_if_else_functor_impl()>> { template @@ -76,8 +88,8 @@ struct copy_if_else_functor_impl auto const& lhs = *p_lhs; auto const& rhs = *p_rhs; - auto lhs_iter = cudf::detail::make_optional_iterator(lhs, nullate::DYNAMIC{left_nullable}); - auto rhs_iter = cudf::detail::make_optional_iterator(rhs, nullate::DYNAMIC{right_nullable}); + auto lhs_iter = make_iterable_optional_iterator(lhs, left_nullable); + auto rhs_iter = make_iterable_optional_iterator(rhs, right_nullable); return detail::copy_if_else(left_nullable || right_nullable, lhs_iter, lhs_iter + size, @@ -111,8 +123,8 @@ struct copy_if_else_functor_impl { auto const& lhs = *p_lhs; auto const& rhs = *p_rhs; - auto lhs_iter = cudf::detail::make_optional_iterator(lhs, nullate::DYNAMIC{left_nullable}); - auto rhs_iter = cudf::detail::make_optional_iterator(rhs, nullate::DYNAMIC{right_nullable}); + auto lhs_iter = make_iterable_optional_iterator(lhs, left_nullable); + auto rhs_iter = make_iterable_optional_iterator(rhs, right_nullable); return strings::detail::copy_if_else(lhs_iter, lhs_iter + size, rhs_iter, filter, stream, mr); } }; diff --git a/cpp/src/copying/get_element.cu b/cpp/src/copying/get_element.cu index 4d9f2867dd76..84bb6e4290dc 100644 --- a/cpp/src/copying/get_element.cu +++ b/cpp/src/copying/get_element.cu @@ -14,7 +14,6 @@ #include #include #include -#include #include #include #include @@ -36,17 +35,16 @@ struct get_element_functor { rmm::device_async_resource_ref mr) { auto s = make_fixed_width_scalar(data_type(type_to_id()), stream, mr); + s->set_valid_async(is_element_valid_sync(input, index, stream), stream); - using ScalarType = cudf::scalar_type_t; - auto typed_s = static_cast(s.get()); - - auto device_s = get_scalar_device_view(*typed_s); + auto output_view = s->as_mutable_column_view(); + auto device_s = + mutable_column_device_view::create(output_view.as_mutable_column_view(), stream); auto device_col = column_device_view::create(input, stream); device_single_thread( - [device_s, d_col = *device_col, index] __device__() mutable { - device_s.set_value(d_col.element(index)); - device_s.set_valid(d_col.is_valid(index)); + [d_scalar = *device_s, d_col = *device_col, index] __device__() mutable { + d_scalar.element(0) = d_col.element(index); }, stream); return s; @@ -84,16 +82,18 @@ struct get_element_functor { { auto dict_view = dictionary_column_view(input); auto indices_iter = detail::indexalator_factory::make_input_iterator(dict_view.indices()); - numeric_scalar key_index_scalar{ - index, true, stream, cudf::get_current_device_resource_ref()}; - auto d_key_index = get_scalar_device_view(key_index_scalar); - auto d_col = column_device_view::create(input, stream); + numeric_scalar key_index_scalar{index, + is_element_valid_sync(input, index, stream), + stream, + cudf::get_current_device_resource_ref()}; + auto key_index_view = key_index_scalar.as_mutable_column_view(); + auto d_key_index = + mutable_column_device_view::create(key_index_view.as_mutable_column_view(), stream); // retrieve the indices value at index device_single_thread( - [d_key_index, d_col = *d_col, indices_iter, index] __device__() mutable { - d_key_index.set_value(indices_iter[index]); - d_key_index.set_valid(d_col.is_valid(index)); + [d_key_index = *d_key_index, indices_iter, index] __device__() mutable { + d_key_index.element(0) = indices_iter[index]; }, stream); @@ -144,16 +144,16 @@ struct get_element_functor { auto device_col = column_device_view::create(input, stream); - auto result = std::make_unique>( - Type{}, numeric::scale_type{input.type().scale()}, false, stream, mr); + auto result = + std::make_unique>(Type{}, + numeric::scale_type{input.type().scale()}, + is_element_valid_sync(input, index, stream), + stream, + mr); device_single_thread( - [buffer = result->data(), - validity = result->validity_data(), - d_col = *device_col, - index] __device__() mutable { - *buffer = d_col.element(index); - *validity = d_col.is_valid(index); + [buffer = result->data(), d_col = *device_col, index] __device__() mutable { + *buffer = d_col.element(index); }, stream); diff --git a/cpp/src/copying/scatter.cu b/cpp/src/copying/scatter.cu index 9fbeabd96135..22df8584597a 100644 --- a/cpp/src/copying/scatter.cu +++ b/cpp/src/copying/scatter.cu @@ -16,7 +16,7 @@ #include #include #include -#include +#include #include #include #include @@ -178,10 +178,10 @@ struct column_scalar_scatterer_impl { cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { + auto const source_view = source.get().as_column_view(); auto dict_target = dictionary::detail::add_keys( dictionary_column_view(target), - make_column_from_scalar(source.get(), 1, stream, cudf::get_current_device_resource_ref()) - ->view(), + source_view.as_column_view(), stream, mr); auto dict_view = dictionary_column_view(dict_target->view()); diff --git a/cpp/src/copying/segmented_shift.cu b/cpp/src/copying/segmented_shift.cu index 99df9ac15dd1..2f4eaff022ee 100644 --- a/cpp/src/copying/segmented_shift.cu +++ b/cpp/src/copying/segmented_shift.cu @@ -65,11 +65,15 @@ struct segmented_shift_functor() rmm::device_async_resource_ref mr) { auto values_device_view = column_device_view::create(segmented_values, stream); - bool nullable = not fill_value.is_valid(stream) or segmented_values.nullable(); + auto const fill_view = fill_value.as_column_view(); + auto fill_device_view = column_device_view::create(fill_view.as_column_view(), stream); + bool nullable = not fill_value.is_valid() or segmented_values.nullable(); auto input_iterator = cudf::detail::make_optional_iterator( *values_device_view, nullate::DYNAMIC{segmented_values.has_nulls()}) - offset; - auto fill_iterator = cudf::detail::make_optional_iterator(fill_value, nullate::YES{}); + auto fill_iterator = cuda::make_permutation_iterator( + cudf::detail::make_optional_iterator(*fill_device_view, nullate::YES{}), + cuda::make_constant_iterator(0)); return copy_if_else(nullable, input_iterator, input_iterator + segmented_values.size(), @@ -94,10 +98,14 @@ struct segmented_shift_functor { rmm::device_async_resource_ref mr) { auto values_device_view = column_device_view::create(segmented_values, stream); + auto const fill_view = fill_value.as_column_view(); + auto fill_device_view = column_device_view::create(fill_view.as_column_view(), stream); auto input_iterator = make_optional_iterator( *values_device_view, nullate::DYNAMIC{segmented_values.has_nulls()}) - offset; - auto fill_iterator = make_optional_iterator(fill_value, nullate::YES{}); + auto fill_iterator = cuda::make_permutation_iterator( + make_optional_iterator(*fill_device_view, nullate::YES{}), + cuda::make_constant_iterator(0)); return strings::detail::copy_if_else(input_iterator, input_iterator + segmented_values.size(), fill_iterator, diff --git a/cpp/src/copying/shift.cu b/cpp/src/copying/shift.cu index 7e7504248340..09be4d684be3 100644 --- a/cpp/src/copying/shift.cu +++ b/cpp/src/copying/shift.cu @@ -40,16 +40,15 @@ inline bool __device__ out_of_bounds(size_type size, size_type idx) std::pair create_null_mask(column_device_view const& input, size_type offset, - scalar const& fill_value, + column_device_view const& fill, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { - auto const size = input.size(); - auto func_validity = - [size, offset, fill = fill_value.validity_data(), input] __device__(size_type idx) { - auto src_idx = idx - offset; - return out_of_bounds(size, src_idx) ? *fill : input.is_valid(src_idx); - }; + auto const size = input.size(); + auto func_validity = [size, offset, fill, input] __device__(size_type idx) { + auto src_idx = idx - offset; + return out_of_bounds(size, src_idx) ? fill.is_valid(0) : input.is_valid(src_idx); + }; return detail::valid_if(cuda::counting_iterator{0}, cuda::counting_iterator{size}, func_validity, @@ -76,9 +75,11 @@ struct shift_functor { auto output = cudf::strings::detail::shift( cudf::strings_column_view(input), offset, fill_value, stream, mr); - if (input.nullable() || not fill_value.is_valid(stream)) { + if (input.nullable() || not fill_value.is_valid()) { auto const d_input = column_device_view::create(input, stream); - auto [null_mask, null_count] = create_null_mask(*d_input, offset, fill_value, stream, mr); + auto const fill_view = fill_value.as_column_view(); + auto const d_fill = column_device_view::create(fill_view.as_column_view(), stream); + auto [null_mask, null_count] = create_null_mask(*d_input, offset, *d_fill, stream, mr); output->set_null_mask(std::move(null_mask), null_count); } @@ -93,19 +94,18 @@ struct shift_functor { rmm::device_async_resource_ref mr) requires(cudf::is_fixed_width()) { - using ScalarType = cudf::scalar_type_t; - auto& scalar = static_cast(fill_value); - - auto device_input = column_device_view::create(input, stream); + auto device_input = column_device_view::create(input, stream); + auto const fill_view = fill_value.as_column_view(); + auto device_fill = column_device_view::create(fill_view.as_column_view(), stream); auto output = detail::allocate_like(input, input.size(), mask_allocation_policy::NEVER, stream, mr); auto device_output = mutable_column_device_view::create(*output, stream); - auto const scalar_is_valid = scalar.is_valid(stream); + auto const scalar_is_valid = fill_value.is_valid(); if (input.nullable() || not scalar_is_valid) { auto [null_mask, null_count] = - create_null_mask(*device_input, offset, fill_value, stream, mr); + create_null_mask(*device_input, offset, *device_fill, stream, mr); output->set_null_mask(std::move(null_mask), null_count); } @@ -126,9 +126,9 @@ struct shift_functor { } auto func_value = - [size, offset, fill = scalar.data(), input = *device_input] __device__(size_type idx) { + [size, offset, fill = *device_fill, input = *device_input] __device__(size_type idx) { auto src_idx = idx - offset; - return out_of_bounds(size, src_idx) ? *fill : input.element(src_idx); + return out_of_bounds(size, src_idx) ? fill.element(0) : input.element(src_idx); }; thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), diff --git a/cpp/src/dictionary/replace.cu b/cpp/src/dictionary/replace.cu index f4697b167909..6b1b5bf99ece 100644 --- a/cpp/src/dictionary/replace.cu +++ b/cpp/src/dictionary/replace.cu @@ -117,12 +117,9 @@ std::unique_ptr replace_nulls(dictionary_column_view const& input, cudf::data_type_error); // first add the replacement to the keys so only the indices need to be processed + auto const replacement_view = replacement.as_column_view(); auto input_matched = dictionary::detail::add_keys( - input, - make_column_from_scalar(replacement, 1, stream, cudf::get_current_device_resource_ref()) - ->view(), - stream, - mr); + input, replacement_view.as_column_view(), stream, mr); auto const input_view = dictionary_column_view(input_matched->view()); auto const scalar_index = get_index(input_view, replacement, stream, cudf::get_current_device_resource_ref()); diff --git a/cpp/src/dictionary/search.cu b/cpp/src/dictionary/search.cu index 3d90c811d443..b02e97c8fb35 100644 --- a/cpp/src/dictionary/search.cu +++ b/cpp/src/dictionary/search.cu @@ -7,7 +7,7 @@ #include #include #include -#include +#include #include #include #include @@ -49,14 +49,16 @@ struct find_index_fn { "search key type must match dictionary keys type", std::invalid_argument); - using ScalarType = cudf::scalar_type_t; - auto const find_key = - get_scalar_device_view(static_cast(const_cast(key))); - auto keys_view = column_device_view::create(input.keys(), stream); - auto const keys = keys_view->begin(); + auto const key_view = key.as_column_view(); + auto find_key_view = column_device_view::create(key_view.as_column_view(), stream); + auto keys_view = column_device_view::create(input.keys(), stream); + auto const find_key = *find_key_view; + auto const keys = keys_view->begin(); auto result = std::make_unique>(-1, true, stream, mr); - auto find_fn = [find_key] __device__(auto const& k) { return k == find_key.value(); }; + auto find_fn = [find_key] __device__(auto const& k) { + return k == find_key.element(0); + }; auto tmp_size = std::size_t{0}; CUDF_CUDA_TRY(cub::DeviceFind::FindIf( nullptr, tmp_size, keys, result->data(), find_fn, num_keys, stream.get())); diff --git a/cpp/src/filling/fill.cu b/cpp/src/filling/fill.cu index 5da1423438a9..f6d4fb0f6261 100644 --- a/cpp/src/filling/fill.cu +++ b/cpp/src/filling/fill.cu @@ -4,7 +4,6 @@ */ #include -#include #include #include #include @@ -155,10 +154,9 @@ std::unique_ptr out_of_place_fill_range_dispatch::operator()view(), stream, mr); + auto const value_view = value.as_column_view(); + auto target_matched = cudf::dictionary::detail::add_keys( + target, value_view.as_column_view(), stream, mr); cudf::column_view const target_indices = cudf::dictionary_column_view(target_matched->view()).get_indices_annotated(); diff --git a/cpp/src/filling/sequence.cu b/cpp/src/filling/sequence.cu index 0af4d908bb48..90e99456e2bb 100644 --- a/cpp/src/filling/sequence.cu +++ b/cpp/src/filling/sequence.cu @@ -10,7 +10,6 @@ #include #include #include -#include #include #include #include @@ -31,20 +30,20 @@ namespace { // __T289 link error. This seems to be related to lambda usage within functions using SFINAE. template struct tabulator { - cudf::numeric_scalar_device_view const n_init; - cudf::numeric_scalar_device_view const n_step; + cudf::column_device_view const init; + cudf::column_device_view const step; T __device__ operator()(cudf::size_type i) { - return n_init.value() + (static_cast(i) * n_step.value()); + return init.element(0) + (static_cast(i) * step.element(0)); } }; template struct const_tabulator { - cudf::numeric_scalar_device_view const n_init; + cudf::column_device_view const init; - T __device__ operator()(cudf::size_type i) { return n_init.value() + static_cast(i); } + T __device__ operator()(cudf::size_type i) { return init.element(0) + static_cast(i); } }; /** @@ -63,10 +62,10 @@ struct sequence_functor { auto result = make_fixed_width_column(init.type(), size, mask_state::UNALLOCATED, stream, mr); auto result_device_view = mutable_column_device_view::create(*result, stream); - auto n_init = - get_scalar_device_view(static_cast&>(const_cast(init))); - auto n_step = - get_scalar_device_view(static_cast&>(const_cast(step))); + auto const init_view = init.as_column_view(); + auto const step_view = step.as_column_view(); + auto const d_init = column_device_view::create(init_view.as_column_view(), stream); + auto const d_step = column_device_view::create(step_view.as_column_view(), stream); // not using thrust::sequence because it requires init and step to be passed as // constants, not iterators. to do that we would have to retrieve the scalar values off the gpu, @@ -74,7 +73,7 @@ struct sequence_functor { thrust::tabulate(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), result_device_view->begin(), result_device_view->end(), - tabulator{n_init, n_step}); + tabulator{*d_init, *d_step}); return result; } @@ -89,8 +88,8 @@ struct sequence_functor { auto result = make_fixed_width_column(init.type(), size, mask_state::UNALLOCATED, stream, mr); auto result_device_view = mutable_column_device_view::create(*result, stream); - auto n_init = - get_scalar_device_view(static_cast&>(const_cast(init))); + auto const init_view = init.as_column_view(); + auto const d_init = column_device_view::create(init_view.as_column_view(), stream); // not using thrust::sequence because it requires init and step to be passed as // constants, not iterators. to do that we would have to retrieve the scalar values off the gpu, @@ -98,7 +97,7 @@ struct sequence_functor { thrust::tabulate(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), result_device_view->begin(), result_device_view->end(), - const_tabulator{n_init}); + const_tabulator{*d_init}); return result; } @@ -125,8 +124,8 @@ std::unique_ptr sequence(size_type size, CUDF_EXPECTS(size >= 0, "size must be >= 0", std::invalid_argument); CUDF_EXPECTS( is_numeric(init.type()), "Input scalar types must be numeric", std::invalid_argument); - CUDF_EXPECTS(init.is_valid(stream), "init must be a valid scalar", std::invalid_argument); - CUDF_EXPECTS(step.is_valid(stream), "step must be a valid scalar", std::invalid_argument); + CUDF_EXPECTS(init.is_valid(), "init must be a valid scalar", std::invalid_argument); + CUDF_EXPECTS(step.is_valid(), "step must be a valid scalar", std::invalid_argument); return type_dispatcher(init.type(), sequence_functor{}, size, init, step, stream, mr); } @@ -138,7 +137,7 @@ std::unique_ptr sequence(size_type size, { CUDF_EXPECTS(size >= 0, "size must be >= 0", std::invalid_argument); CUDF_EXPECTS(is_numeric(init.type()), "init scalar type must be numeric", cudf::data_type_error); - CUDF_EXPECTS(init.is_valid(stream), "init must be a valid scalar", std::invalid_argument); + CUDF_EXPECTS(init.is_valid(), "init must be a valid scalar", std::invalid_argument); return type_dispatcher(init.type(), sequence_functor{}, size, init, stream, mr); } diff --git a/cpp/src/groupby/hash/compute_single_pass_aggs.cuh b/cpp/src/groupby/hash/compute_single_pass_aggs.cuh index fd248138fabf..b10718a8fa63 100644 --- a/cpp/src/groupby/hash/compute_single_pass_aggs.cuh +++ b/cpp/src/groupby/hash/compute_single_pass_aggs.cuh @@ -18,6 +18,7 @@ #include #include +#include #include #include diff --git a/cpp/src/io/utilities/data_casting.cu b/cpp/src/io/utilities/data_casting.cu index 111048e8e9fe..ead99544193d 100644 --- a/cpp/src/io/utilities/data_casting.cu +++ b/cpp/src/io/utilities/data_casting.cu @@ -14,6 +14,7 @@ #include #include #include +#include #include #include #include diff --git a/cpp/src/jit/row_ir.cpp b/cpp/src/jit/row_ir.cpp index b39a1d6f0e7b..e4ae251d0184 100644 --- a/cpp/src/jit/row_ir.cpp +++ b/cpp/src/jit/row_ir.cpp @@ -822,7 +822,7 @@ if(expected__{1}.has_value()) {{ std::unique_ptr ast_converter::add_ir_node(ast::literal const& expr) { auto id = expr.is_scalar_column_view() ? instance_.add_input(expr.get_scalar_column_view()) - : instance_.add_input(expr.get_scalar()); + : instance_.add_input(expr.get_scalar().as_column_view()); return std::make_unique(input_reference{id}); } @@ -876,9 +876,9 @@ std::unique_ptr ast_converter::add_ir_node(ast::jit::detail::opera bool is_nullable(scalar_input const& in) { if (auto* s = std::get_if>(&in)) { - return (*s)->nullable(); + return (*s)->has_nulls(); } else { - return std::get(in).nullable(); + return std::get(in).has_nulls(); } } diff --git a/cpp/src/reductions/simple.cuh b/cpp/src/reductions/simple.cuh index 319613e0622b..fd8b6c3c1727 100644 --- a/cpp/src/reductions/simple.cuh +++ b/cpp/src/reductions/simple.cuh @@ -7,6 +7,7 @@ #include "nested_types_extrema_utils.cuh" +#include #include #include #include @@ -14,7 +15,6 @@ #include #include #include -#include #include #include #include @@ -176,12 +176,11 @@ template struct assign_scalar_fn { __device__ void operator()() { - d_output.set_value(static_cast(d_input.value())); - d_output.set_valid(d_input.is_valid()); + d_output.element(0) = static_cast(d_input.element(0)); } - cudf::numeric_scalar_device_view d_input; - cudf::numeric_scalar_device_view d_output; + cudf::column_device_view d_input; + cudf::mutable_column_device_view d_output; }; /** @@ -208,10 +207,14 @@ struct cast_numeric_scalar_fn { rmm::device_async_resource_ref mr) requires(is_supported()) { - auto d_input = cudf::get_scalar_device_view(*input); - auto result = std::make_unique>(ResultType{}, true, stream, mr); - auto d_output = cudf::get_scalar_device_view(*result); - cudf::detail::device_single_thread(assign_scalar_fn{d_input, d_output}, + auto result = + std::make_unique>(ResultType{}, input->is_valid(), stream, mr); + auto const input_view = input->as_column_view(); + auto output_view = result->as_mutable_column_view(); + auto d_input = cudf::column_device_view::create(input_view.as_column_view(), stream); + auto d_output = + cudf::mutable_column_device_view::create(output_view.as_mutable_column_view(), stream); + cudf::detail::device_single_thread(assign_scalar_fn{*d_input, *d_output}, stream); return result; } diff --git a/cpp/src/replace/clamp.cu b/cpp/src/replace/clamp.cu index 836af9b5a86a..87d41ec6a921 100644 --- a/cpp/src/replace/clamp.cu +++ b/cpp/src/replace/clamp.cu @@ -17,7 +17,6 @@ #include #include #include -#include #include #include #include @@ -130,12 +129,9 @@ std::unique_ptr clamp_dictionary_column(dictionary_column_view con std::unique_ptr result = nullptr; auto add_scalar_key = [&](scalar const& key, scalar const& key_replace) { if (key.is_valid(stream)) { + auto const key_replace_view = key_replace.as_column_view(); result = dictionary::detail::add_keys( - matched_view, - make_column_from_scalar(key_replace, 1, stream, cudf::get_current_device_resource_ref()) - ->view(), - stream, - mr); + matched_view, key_replace_view.as_column_view(), stream, mr); matched_view = dictionary_column_view(result->view()); } }; @@ -159,8 +155,14 @@ std::unique_ptr clamp_dictionary_column(dictionary_column_view con auto indices_itr = cudf::detail::indexalator_factory::make_output_iterator(indices_column->mutable_view()); - auto lo_itr = make_optional_iterator(lo, nullate::YES{}); - auto hi_itr = make_optional_iterator(hi, nullate::YES{}); + auto const lo_view = lo.as_column_view(); + auto const hi_view = hi.as_column_view(); + auto d_lo = column_device_view::create(lo_view.as_column_view(), stream); + auto d_hi = column_device_view::create(hi_view.as_column_view(), stream); + auto lo_itr = cuda::make_permutation_iterator(make_optional_iterator(*d_lo, nullate::YES{}), + cuda::make_constant_iterator(0)); + auto hi_itr = cuda::make_permutation_iterator(make_optional_iterator(*d_hi, nullate::YES{}), + cuda::make_constant_iterator(0)); auto d_input = column_device_view::create(input.parent(), stream); using OptionalIterator = decltype(lo_itr); @@ -280,10 +282,24 @@ struct dispatch_clamp { dictionary_column_view(input), lo, lo_replace, hi, hi_replace, stream, mr); } - auto lo_itr = make_optional_iterator(lo, nullate::YES{}); - auto hi_itr = make_optional_iterator(hi, nullate::YES{}); - auto lo_replace_itr = make_optional_iterator(lo_replace, nullate::NO{}); - auto hi_replace_itr = make_optional_iterator(hi_replace, nullate::NO{}); + auto const lo_view = lo.as_column_view(); + auto const hi_view = hi.as_column_view(); + auto const lo_replace_view = lo_replace.as_column_view(); + auto const hi_replace_view = hi_replace.as_column_view(); + auto d_lo = column_device_view::create(lo_view.as_column_view(), stream); + auto d_hi = column_device_view::create(hi_view.as_column_view(), stream); + auto d_lo_replace = column_device_view::create(lo_replace_view.as_column_view(), stream); + auto d_hi_replace = column_device_view::create(hi_replace_view.as_column_view(), stream); + + auto const scalar_index = cuda::make_constant_iterator(0); + auto lo_itr = cuda::make_permutation_iterator(make_optional_iterator(*d_lo, nullate::YES{}), + scalar_index); + auto hi_itr = cuda::make_permutation_iterator(make_optional_iterator(*d_hi, nullate::YES{}), + scalar_index); + auto lo_replace_itr = cuda::make_permutation_iterator( + make_optional_iterator(*d_lo_replace, nullate::NO{}), scalar_index); + auto hi_replace_itr = cuda::make_permutation_iterator( + make_optional_iterator(*d_hi_replace, nullate::NO{}), scalar_index); return clamp(input, lo_itr, lo_replace_itr, hi_itr, hi_replace_itr, stream, mr); } diff --git a/cpp/src/replace/nans.cu b/cpp/src/replace/nans.cu index c943b7a78be8..44bd7fe0affc 100644 --- a/cpp/src/replace/nans.cu +++ b/cpp/src/replace/nans.cu @@ -33,6 +33,7 @@ struct replace_nans_functor { std::unique_ptr operator()(column_view const& input, Replacement const& replacement, bool replacement_nullable, + bool broadcast_replacement, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) requires(std::is_floating_point_v) @@ -51,8 +52,14 @@ struct replace_nans_functor { auto input_iterator = make_optional_iterator(*input_device_view, nullate::DYNAMIC{input.has_nulls()}); - auto replacement_iterator = + auto replacement_values = make_optional_iterator(replacement, nullate::DYNAMIC{replacement_nullable}); + auto replacement_indices = make_counting_transform_iterator( + size_type{0}, [broadcast_replacement] __device__(size_type i) -> size_type { + return broadcast_replacement ? 0 : i; + }); + auto replacement_iterator = + cuda::make_permutation_iterator(replacement_values, replacement_indices); return copy_if_else(input.has_nulls() or replacement_nullable, input_iterator, input_iterator + size, @@ -81,11 +88,13 @@ std::unique_ptr replace_nans(column_view const& input, CUDF_EXPECTS(input.size() == replacement.size(), "Input and replacement must be of the same size"); + auto replacement_device_view = column_device_view::create(replacement, stream); return type_dispatcher(input.type(), replace_nans_functor{}, input, - *column_device_view::create(replacement, stream), + *replacement_device_view, replacement.nullable(), + false, stream, mr); } @@ -95,8 +104,11 @@ std::unique_ptr replace_nans(column_view const& input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { + auto const replacement_view = replacement.as_column_view(); + auto replacement_device_view = + column_device_view::create(replacement_view.as_column_view(), stream); return type_dispatcher( - input.type(), replace_nans_functor{}, input, replacement, true, stream, mr); + input.type(), replace_nans_functor{}, input, *replacement_device_view, true, true, stream, mr); } } // namespace detail diff --git a/cpp/src/scalar/scalar.cpp b/cpp/src/scalar/scalar.cpp index 9fe9befb4ea0..5de16a85d4e5 100644 --- a/cpp/src/scalar/scalar.cpp +++ b/cpp/src/scalar/scalar.cpp @@ -4,6 +4,7 @@ */ #include +#include #include #include #include @@ -22,58 +23,219 @@ namespace cudf { -static rmm::device_buffer make_string_device_buffer(std::string_view string, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) +namespace { + +rmm::device_buffer make_null_mask(bool is_valid, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + return cudf::detail::create_null_mask( + 1, is_valid ? mask_state::ALL_VALID : mask_state::ALL_NULL, stream, mr); +} + +size_type checked_string_size(std::size_t size) +{ + CUDF_EXPECTS(size <= static_cast(std::numeric_limits::max()), + "Data exceeds the string size limit", + std::overflow_error); + return static_cast(size); +} + +template +data_type fixed_width_storage_type() +{ + if constexpr (std::is_same_v) { + return data_type{type_id::DECIMAL128, 0}; + } else { + return data_type{type_to_id()}; + } +} + +template +column make_fixed_width_storage(T value, + data_type type, + bool is_valid, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + auto host_data = cudf::detail::make_pinned_vector(1, stream); + host_data[0] = value; + return column{type, + 1, + rmm::device_buffer{host_data.data(), sizeof(T), stream, mr}, + make_null_mask(is_valid, stream, mr), + is_valid ? 0 : 1}; +} + +template +column make_fixed_width_storage(rmm::device_scalar const& value, + data_type type, + bool is_valid, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + rmm::device_buffer data(sizeof(T), stream, mr); + CUDF_CUDA_TRY( + cudaMemcpyAsync(data.data(), value.data(), sizeof(T), cudaMemcpyDeviceToDevice, stream.get())); + return column{type, 1, std::move(data), make_null_mask(is_valid, stream, mr), is_valid ? 0 : 1}; +} + +std::unique_ptr make_offsets_column(size_type size, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + auto offsets = cudf::detail::make_pinned_vector(2, stream); + offsets[0] = 0; + offsets[1] = size; + return std::make_unique( + data_type{type_id::INT32}, + 2, + rmm::device_buffer{offsets.data(), offsets.size() * sizeof(size_type), stream, mr}, + rmm::device_buffer{}, + 0); +} + +column assemble_string_storage(rmm::device_buffer&& chars, + std::unique_ptr&& offsets, + bool is_valid, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + std::vector> children; + children.push_back(std::move(offsets)); + return column{data_type{type_id::STRING}, + 1, + std::move(chars), + make_null_mask(is_valid, stream, mr), + is_valid ? 0 : 1, + std::move(children)}; +} + +column make_string_storage(rmm::device_buffer&& chars, + size_type size, + bool is_valid, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + return assemble_string_storage( + std::move(chars), make_offsets_column(size, stream, mr), is_valid, stream, mr); +} + +column make_list_storage(column&& elements, + bool is_valid, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + auto const size = elements.size(); + std::vector> children; + children.push_back(make_offsets_column(size, stream, mr)); + children.push_back(std::make_unique(std::move(elements))); + return column{data_type{type_id::LIST}, + 1, + rmm::device_buffer{}, + make_null_mask(is_valid, stream, mr), + is_valid ? 0 : 1, + std::move(children)}; +} + +column make_struct_storage(table&& children_table, + bool is_valid, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + return column{data_type{type_id::STRUCT}, + 1, + rmm::device_buffer{}, + make_null_mask(is_valid, stream, mr), + is_valid ? 0 : 1, + children_table.release()}; +} + +} // namespace + +struct string_scalar::string_storage { + column storage; + cudf::detail::host_vector staging; + size_type size; +}; + +string_scalar::string_storage string_scalar::make_storage(std::string_view string, + bool is_valid, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + auto const size = checked_string_size(string.size()); + auto offsets = make_offsets_column(size, stream, mr); + auto staging = cudf::detail::make_pinned_vector(string.size(), stream); + std::copy(string.begin(), string.end(), staging.begin()); + auto chars = rmm::device_buffer(staging.data(), staging.size(), stream, mr); + return string_storage{ + assemble_string_storage(std::move(chars), std::move(offsets), is_valid, stream, mr), + std::move(staging), + size}; +} + +string_scalar::string_scalar(string_storage&& storage) + : scalar(std::move(storage.storage)), _size(storage.size), _host_data(std::move(storage.staging)) { - auto host_data = cudf::detail::make_pinned_vector(string.size(), stream); - std::copy(string.begin(), string.end(), host_data.begin()); - return rmm::device_buffer(host_data.data(), host_data.size(), stream, mr); } scalar::scalar(data_type type, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : _type(type), _is_valid(is_valid, stream, mr) + : _storage{type, + 1, + rmm::device_buffer{cudf::size_of(type), stream, mr}, + make_null_mask(is_valid, stream, mr), + is_valid ? 0 : 1} { } +scalar::scalar(column&& storage) : _storage(std::move(storage)) +{ + CUDF_EXPECTS(_storage.size() == 1, "Scalar storage must contain exactly one row."); +} + scalar::scalar(scalar const& other, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : _type(other.type()), _is_valid(other._is_valid, stream, mr) + : _storage(other._storage, stream, mr) { } -data_type scalar::type() const noexcept { return _type; } +data_type scalar::type() const noexcept { return _storage.type(); } void scalar::set_valid_async(bool is_valid, cuda::stream_ref stream) { - _is_valid.set_value_async(is_valid, stream); + CUDF_CUDA_TRY(cudaMemsetAsync(_storage.mutable_view().null_mask(), + is_valid ? 0xff : 0, + bitmask_allocation_size_bytes(1), + stream.get())); + _storage.set_null_count(is_valid ? 0 : 1); } -bool scalar::is_valid(cuda::stream_ref stream) const { return _is_valid.value(stream); } +bool scalar::is_valid(cuda::stream_ref) const { return _storage.null_count() == 0; } -bool* scalar::validity_data() { return _is_valid.data(); } +bitmask_type const* scalar::validity_data() const { return _storage.view().null_mask(); } -bool const* scalar::validity_data() const { return _is_valid.data(); } +scalar_column_view scalar::as_column_view() const { return scalar_column_view{_storage.view()}; } + +mutable_scalar_column_view scalar::as_mutable_column_view() +{ + return mutable_scalar_column_view{_storage.mutable_view()}; +} string_scalar::string_scalar(std::string_view string, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar(data_type(type_id::STRING), is_valid, stream, mr), - _data(make_string_device_buffer(string, stream, mr)) + : string_scalar(make_storage(string, is_valid, stream, mr)) { - CUDF_EXPECTS( - string.size() <= static_cast(std::numeric_limits::max()), - "Data exceeds the string size limit", - std::overflow_error); } string_scalar::string_scalar(string_scalar const& other, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar(other, stream, mr), _data(other._data, stream, mr) + : scalar(other, stream, mr), _size(other._size) { } @@ -89,16 +251,30 @@ string_scalar::string_scalar(value_type const& source, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar(data_type(type_id::STRING), is_valid, stream, mr), - _data(source.data(), source.size_bytes(), stream, mr) + : scalar([&] { + rmm::device_buffer chars(source.size_bytes(), stream, mr); + CUDF_CUDA_TRY(cudaMemcpyAsync( + chars.data(), source.data(), source.size_bytes(), cudaMemcpyDeviceToDevice, stream.get())); + return make_string_storage(std::move(chars), source.size_bytes(), is_valid, stream, mr); + }()), + _size(source.size_bytes()) +{ +} + +string_scalar::string_scalar(rmm::device_buffer&& data, + bool is_valid, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) + : string_scalar(std::move(data), checked_string_size(data.size()), is_valid, stream, mr) { } string_scalar::string_scalar(rmm::device_buffer&& data, + size_type size, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar(data_type(type_id::STRING), is_valid, stream, mr), _data(std::move(data)) + : scalar(make_string_storage(std::move(data), size, is_valid, stream, mr)), _size(size) { } @@ -107,15 +283,15 @@ string_scalar::value_type string_scalar::value(cuda::stream_ref stream) const return value_type{data(), size()}; } -size_type string_scalar::size() const { return _data.size(); } +size_type string_scalar::size() const { return _size; } -char const* string_scalar::data() const { return static_cast(_data.data()); } +char const* string_scalar::data() const { return _storage.view().data(); } std::string string_scalar::to_string(cuda::stream_ref stream) const { std::string result(size(), '\0'); detail::cuda_memcpy(host_span{result.data(), result.size()}, - device_span{data(), _data.size()}, + device_span{data(), static_cast(size())}, stream); return result; } @@ -126,8 +302,8 @@ fixed_point_scalar::fixed_point_scalar(rep_type value, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar{data_type{type_to_id(), static_cast(scale)}, is_valid, stream, mr}, - _data{value, stream, mr} + : scalar{make_fixed_width_storage( + value, data_type{type_to_id(), static_cast(scale)}, is_valid, stream, mr)} { } @@ -136,7 +312,7 @@ fixed_point_scalar::fixed_point_scalar(rep_type value, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar{data_type{type_to_id(), 0}, is_valid, stream, mr}, _data{value, stream, mr} + : scalar{make_fixed_width_storage(value, data_type{type_to_id(), 0}, is_valid, stream, mr)} { } @@ -145,8 +321,8 @@ fixed_point_scalar::fixed_point_scalar(T value, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar{data_type{type_to_id(), value.scale()}, is_valid, stream, mr}, - _data{value.value(), stream, mr} + : scalar{make_fixed_width_storage( + value.value(), data_type{type_to_id(), value.scale()}, is_valid, stream, mr)} { } @@ -156,8 +332,7 @@ fixed_point_scalar::fixed_point_scalar(rmm::device_scalar&& data, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar{data_type{type_to_id(), scale}, is_valid, stream, mr}, - _data{data.value(stream), stream, mr} + : scalar{make_fixed_width_storage(data, data_type{type_to_id(), scale}, is_valid, stream, mr)} { } @@ -165,33 +340,36 @@ template fixed_point_scalar::fixed_point_scalar(fixed_point_scalar const& other, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar{other, stream, mr}, _data(other._data, stream, mr) + : scalar{other, stream, mr} { } template typename fixed_point_scalar::rep_type fixed_point_scalar::value(cuda::stream_ref stream) const { - return _data.value(stream); + rep_type result; + detail::cuda_memcpy( + host_span{&result, 1}, device_span{data(), 1}, stream); + return result; } template T fixed_point_scalar::fixed_point_value(cuda::stream_ref stream) const { return value_type{ - numeric::scaled_integer{_data.value(stream), numeric::scale_type{type().scale()}}}; + numeric::scaled_integer{value(stream), numeric::scale_type{type().scale()}}}; } template typename fixed_point_scalar::rep_type* fixed_point_scalar::data() { - return _data.data(); + return _storage.mutable_view().data(); } template typename fixed_point_scalar::rep_type const* fixed_point_scalar::data() const { - return _data.data(); + return _storage.view().data(); } /** @@ -213,7 +391,8 @@ fixed_width_scalar::fixed_width_scalar(T value, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar(data_type(type_to_id()), is_valid, stream, mr), _data(value, stream, mr) + : scalar(make_fixed_width_storage(value, fixed_width_storage_type(), is_valid, stream, mr)), + _bounce_buffer(cudf::detail::make_pinned_vector(1, stream)) { } @@ -222,7 +401,8 @@ fixed_width_scalar::fixed_width_scalar(rmm::device_scalar&& data, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar(data_type(type_to_id()), is_valid, stream, mr), _data{data.value(stream), stream, mr} + : scalar(make_fixed_width_storage(data, fixed_width_storage_type(), is_valid, stream, mr)), + _bounce_buffer(cudf::detail::make_pinned_vector(1, stream)) { } @@ -230,33 +410,36 @@ template fixed_width_scalar::fixed_width_scalar(fixed_width_scalar const& other, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar{other, stream, mr}, _data(other._data, stream, mr) + : scalar{other, stream, mr}, _bounce_buffer(cudf::detail::make_pinned_vector(1, stream)) { } template void fixed_width_scalar::set_value(T value, cuda::stream_ref stream) { - _data.set_value_async(value, stream); + _bounce_buffer[0] = value; + detail::cuda_memcpy_async(device_span{data(), 1}, _bounce_buffer, stream); this->set_valid_async(true, stream); } template T fixed_width_scalar::value(cuda::stream_ref stream) const { - return _data.value(stream); + T result; + detail::cuda_memcpy(host_span{&result, 1}, device_span{data(), 1}, stream); + return result; } template T* fixed_width_scalar::data() { - return _data.data(); + return _storage.mutable_view().data(); } template T const* fixed_width_scalar::data() const { - return _data.data(); + return _storage.view().data(); } /** @@ -492,7 +675,7 @@ list_scalar::list_scalar(cudf::column_view const& data, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar(data_type(type_id::LIST), is_valid, stream, mr), _data(data, stream, mr) + : scalar(make_list_storage(column{data, stream, mr}, is_valid, stream, mr)) { } @@ -500,23 +683,26 @@ list_scalar::list_scalar(cudf::column&& data, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar(data_type(type_id::LIST), is_valid, stream, mr), _data(std::move(data)) + : scalar(make_list_storage(std::move(data), is_valid, stream, mr)) { } list_scalar::list_scalar(list_scalar const& other, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar{other, stream, mr}, _data(other._data, stream, mr) + : scalar{other, stream, mr} { } -column_view list_scalar::view() const { return _data.view(); } +column_view list_scalar::view() const +{ + return _storage.num_children() == 0 ? column_view{} : _storage.view().child(1); +} struct_scalar::struct_scalar(struct_scalar const& other, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar{other, stream, mr}, _data(other._data, stream, mr) + : scalar{other, stream, mr} { } @@ -524,8 +710,8 @@ struct_scalar::struct_scalar(table_view const& data, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar(data_type(type_id::STRUCT), is_valid, stream, mr), - _data{init_data(table{data, stream, mr}, is_valid, stream, mr)} + : scalar(make_struct_storage( + init_data(table{data, stream, mr}, is_valid, stream, mr), is_valid, stream, mr)) { assert_valid_size(); } @@ -534,12 +720,14 @@ struct_scalar::struct_scalar(std::span data, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar(data_type(type_id::STRUCT), is_valid, stream, mr), - _data{ + : scalar(make_struct_storage( init_data(table{table_view{std::vector{data.begin(), data.end()}}, stream, mr}, is_valid, stream, - mr)} + mr), + is_valid, + stream, + mr)) { assert_valid_size(); } @@ -548,17 +736,21 @@ struct_scalar::struct_scalar(table&& data, bool is_valid, cuda::stream_ref stream, rmm::device_async_resource_ref mr) - : scalar(data_type(type_id::STRUCT), is_valid, stream, mr), - _data{init_data(std::move(data), is_valid, stream, mr)} + : scalar( + make_struct_storage(init_data(std::move(data), is_valid, stream, mr), is_valid, stream, mr)) { assert_valid_size(); } -table_view struct_scalar::view() const { return _data.view(); } +table_view struct_scalar::view() const +{ + auto const storage = _storage.view(); + return table_view{std::vector{storage.child_begin(), storage.child_end()}}; +} void struct_scalar::assert_valid_size() { - auto const tv = _data.view(); + auto const tv = view(); CUDF_EXPECTS( std::all_of(tv.begin(), tv.end(), [](column_view const& col) { return col.size() == 1; }), "Struct scalar inputs must have exactly 1 row"); diff --git a/cpp/src/search/contains_scalar.cu b/cpp/src/search/contains_scalar.cu index 1dc60a4aad4a..ea0227dc1a5c 100644 --- a/cpp/src/search/contains_scalar.cu +++ b/cpp/src/search/contains_scalar.cu @@ -3,7 +3,6 @@ * SPDX-License-Identifier: Apache-2.0 */ -#include #include #include #include @@ -12,7 +11,6 @@ #include #include #include -#include #include #include #include @@ -30,23 +28,6 @@ namespace detail { namespace { -/** - * @brief Get the underlying value of a scalar through a scalar device view. - * - * @tparam Element The scalar's value type - * @tparam ScalarDView Type of the input scalar device view - * @param d_scalar The input scalar device view - */ -template -__device__ auto inline get_scalar_value(ScalarDView d_scalar) -{ - if constexpr (cudf::is_fixed_point()) { - return d_scalar.rep(); - } else { - return d_scalar.value(); - } -} - 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. @@ -61,10 +42,10 @@ struct contains_scalar_dispatch { // Don't need to check for needle validity. If it is invalid, it should be handled by the caller // before dispatching to this function. - using DType = device_storage_type_t; - auto const d_haystack = column_device_view::create(haystack, stream); - auto const d_needle = get_scalar_device_view( - static_cast&>(const_cast(needle))); + using DType = device_storage_type_t; + auto const d_haystack = column_device_view::create(haystack, stream); + auto const needle_view = needle.as_column_view(); + auto const d_needle = column_device_view::create(needle_view.as_column_view(), stream); auto const begin = d_haystack->optional_begin(cudf::nullate::DYNAMIC{haystack.has_nulls()}); @@ -73,9 +54,8 @@ struct contains_scalar_dispatch { return cudf::detail::count_if( begin, end, - [d_needle] __device__(auto const val_pair) { - auto needle = get_scalar_value(d_needle); - return val_pair.has_value() && (needle == *val_pair); + [d_needle = *d_needle] __device__(auto const val_pair) { + return val_pair.has_value() && (d_needle.element(0) == *val_pair); }, stream) > 0; } @@ -93,11 +73,11 @@ struct contains_scalar_dispatch { // In addition, haystack and needle structure compatibility will be checked later on by // constructor of the table comparator. - auto const haystack_tv = table_view{{haystack}}; - auto const temp_mr = cudf::get_current_device_resource_ref(); - auto const needle_as_col = make_column_from_scalar(needle, 1, stream, temp_mr); - auto const needle_tv = table_view{{needle_as_col->view()}}; - auto const has_nulls = has_nested_nulls(haystack_tv) || has_nested_nulls(needle_tv); + auto const haystack_tv = table_view{{haystack}}; + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const needle_view = needle.as_column_view(); + auto const needle_tv = table_view{{needle_view.as_column_view()}}; + auto const has_nulls = has_nested_nulls(haystack_tv) || has_nested_nulls(needle_tv); auto const comparator = cudf::detail::row::equality::two_table_comparator(haystack_tv, needle_tv, stream, temp_mr); @@ -143,11 +123,11 @@ bool contains_scalar_dispatch::operator()(column_view const& auto const index = cudf::dictionary::detail::get_index( dict_col, needle, stream, cudf::get_current_device_resource_ref()); // if found, check the index is actually in the indices column - return index->is_valid(stream) && cudf::type_dispatcher(dict_col.indices().type(), - contains_scalar_dispatch{}, - dict_col.indices(), - *index, - stream); + return index->is_valid() && cudf::type_dispatcher(dict_col.indices().type(), + contains_scalar_dispatch{}, + dict_col.indices(), + *index, + stream); } } // namespace @@ -155,7 +135,7 @@ bool contains_scalar_dispatch::operator()(column_view const& bool contains(column_view const& haystack, scalar const& needle, cuda::stream_ref stream) { if (haystack.is_empty()) { return false; } - if (not needle.is_valid(stream)) { return haystack.has_nulls(); } + if (not needle.is_valid()) { return haystack.has_nulls(); } return cudf::type_dispatcher( haystack.type(), contains_scalar_dispatch{}, haystack, needle, stream); diff --git a/cpp/src/strings/combine/concatenate.cu b/cpp/src/strings/combine/concatenate.cu index d3fedcdfc596..c37dff82ed96 100644 --- a/cpp/src/strings/combine/concatenate.cu +++ b/cpp/src/strings/combine/concatenate.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include #include #include @@ -37,7 +37,7 @@ namespace { struct concat_strings_base { table_device_view const d_table; - string_scalar_device_view const d_narep; + column_device_view const d_narep; separator_on_nulls separate_nulls; size_type* d_sizes; char* d_chars; @@ -55,7 +55,7 @@ struct concat_strings_base { */ __device__ void process_row(size_type idx, string_view const d_separator) { - if (!d_narep.is_valid() && + if (!d_narep.is_valid(0) && thrust::any_of(thrust::seq, d_table.begin(), d_table.end(), [idx](auto const& col) { return col.is_null(idx); })) { @@ -78,7 +78,8 @@ struct concat_strings_base { } // write out column's row data (or narep if the row is null) - auto const d_str = null_element ? d_narep.value() : d_column.element(idx); + auto const d_str = + null_element ? d_narep.element(0) : d_column.element(idx); if (d_buffer) d_buffer = detail::copy_string(d_buffer, d_str); bytes += d_str.size_bytes(); @@ -94,17 +95,20 @@ struct concat_strings_base { * @brief Single separator concatenate functor */ struct concat_strings_fn : concat_strings_base { - string_view const d_separator; + column_device_view const d_separator; concat_strings_fn(table_device_view const& d_table, - string_view const& d_separator, - string_scalar_device_view const& d_narep, + column_device_view const& d_separator, + column_device_view const& d_narep, separator_on_nulls separate_nulls) : concat_strings_base{d_table, d_narep, separate_nulls}, d_separator(d_separator) { } - __device__ void operator()(std::size_t idx) { process_row(idx, d_separator); } + __device__ void operator()(std::size_t idx) + { + process_row(idx, d_separator.element(0)); + } }; } // namespace @@ -127,21 +131,23 @@ std::unique_ptr concatenate(table_view const& strings_columns, if (strings_count == 0) // empty begets empty return make_empty_column(type_id::STRING); - CUDF_EXPECTS(separator.is_valid(stream), "Parameter separator must be a valid string_scalar"); - string_view d_separator(separator.data(), separator.size()); - auto d_narep = get_scalar_device_view(const_cast(narep)); + CUDF_EXPECTS(separator.is_valid(), "Parameter separator must be a valid string_scalar"); + auto const separator_view = separator.as_column_view(); + auto const narep_view = narep.as_column_view(); + auto const d_separator = column_device_view::create(separator_view.as_column_view(), stream); + auto const d_narep = column_device_view::create(narep_view.as_column_view(), stream); // Create device views from the strings columns. auto d_table = table_device_view::create(strings_columns, stream); - concat_strings_fn fn{*d_table, d_separator, d_narep, separate_nulls}; + concat_strings_fn fn{*d_table, *d_separator, *d_narep, separate_nulls}; auto [offsets_column, chars] = make_strings_children(fn, strings_count, stream, mr); // create resulting null mask auto [null_mask, null_count] = cudf::detail::valid_if( cuda::counting_iterator{0}, cuda::counting_iterator{strings_count}, - [d_table = *d_table, d_narep] __device__(size_type idx) { - if (d_narep.is_valid()) return true; + [d_table = *d_table, d_narep = *d_narep] __device__(size_type idx) { + if (d_narep.is_valid(0)) return true; return !thrust::any_of( thrust::seq, d_table.begin(), d_table.end(), [idx](auto col) { return col.is_null(idx); }); }, @@ -163,12 +169,12 @@ namespace { */ struct multi_separator_concat_fn : concat_strings_base { column_device_view const d_separators; - string_scalar_device_view const d_separator_narep; + column_device_view const d_separator_narep; multi_separator_concat_fn(table_device_view const& d_table, column_device_view const& d_separators, - string_scalar_device_view const& d_separator_narep, - string_scalar_device_view const& d_narep, + column_device_view const& d_separator_narep, + column_device_view const& d_narep, separator_on_nulls separate_nulls) : concat_strings_base{d_table, d_narep, separate_nulls}, d_separators(d_separators), @@ -178,13 +184,13 @@ struct multi_separator_concat_fn : concat_strings_base { __device__ void operator()(size_type idx) { - if (d_separators.is_null(idx) && !d_separator_narep.is_valid()) { + if (d_separators.is_null(idx) && !d_separator_narep.is_valid(0)) { if (!d_chars) { d_sizes[idx] = 0; } return; } auto const d_separator = d_separators.is_valid(idx) ? d_separators.element(idx) - : d_separator_narep.value(); + : d_separator_narep.element(0); // base class utility function handles the rest process_row(idx, d_separator); } @@ -214,10 +220,11 @@ std::unique_ptr concatenate(table_view const& strings_columns, if (strings_count == 0) // Empty begets empty return make_empty_column(type_id::STRING); - // Invalid output column strings - null rows - string_view const invalid_str{nullptr, 0}; - auto const separator_rep = get_scalar_device_view(const_cast(separator_narep)); - auto const col_rep = get_scalar_device_view(const_cast(col_narep)); + auto const separator_narep_view = separator_narep.as_column_view(); + auto const col_narep_view = col_narep.as_column_view(); + auto const separator_rep = + column_device_view::create(separator_narep_view.as_column_view(), stream); + auto const col_rep = column_device_view::create(col_narep_view.as_column_view(), stream); auto const separator_col_view_ptr = column_device_view::create(separators.parent(), stream); auto const separator_col_view = *separator_col_view_ptr; @@ -225,16 +232,19 @@ std::unique_ptr concatenate(table_view const& strings_columns, auto d_table = table_device_view::create(strings_columns, stream); multi_separator_concat_fn mscf{ - *d_table, separator_col_view, separator_rep, col_rep, separate_nulls}; + *d_table, separator_col_view, *separator_rep, *col_rep, separate_nulls}; auto [offsets_column, chars] = make_strings_children(mscf, strings_count, stream, mr); // Create resulting null mask auto [null_mask, null_count] = cudf::detail::valid_if( cuda::counting_iterator{0}, cuda::counting_iterator{strings_count}, - [d_table = *d_table, separator_col_view, separator_rep, col_rep] __device__(size_type idx) { - if (!separator_col_view.is_valid(idx) && !separator_rep.is_valid()) return false; - if (col_rep.is_valid()) return true; + [d_table = *d_table, + separator_col_view, + separator_rep = *separator_rep, + col_rep = *col_rep] __device__(size_type idx) { + if (!separator_col_view.is_valid(idx) && !separator_rep.is_valid(0)) return false; + if (col_rep.is_valid(0)) return true; return !thrust::any_of( thrust::seq, d_table.begin(), d_table.end(), [idx](auto col) { return col.is_null(idx); }); }, diff --git a/cpp/src/strings/combine/join.cu b/cpp/src/strings/combine/join.cu index 7d0ad215743f..04a4f1a7fb7a 100644 --- a/cpp/src/strings/combine/join.cu +++ b/cpp/src/strings/combine/join.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include #include #include @@ -48,16 +48,16 @@ constexpr size_type AVG_CHAR_BYTES_THRESHOLD = 32; struct join_base_fn { column_device_view const d_strings; - string_view d_separator; - string_scalar_device_view d_narep; + column_device_view const d_separator; + column_device_view const d_narep; __device__ cuda::std::pair process_string(size_type idx) const { string_view d_str{}; - string_view d_sep = (idx + 1 < d_strings.size()) ? d_separator : d_str; + string_view d_sep = (idx + 1 < d_strings.size()) ? d_separator.element(0) : d_str; if (d_strings.is_null(idx)) { - if (d_narep.is_valid()) { - d_str = d_narep.value(); + if (d_narep.is_valid(0)) { + d_str = d_narep.element(0); } else { // if null and no narep, don't output a separator either d_sep = d_str; @@ -80,8 +80,8 @@ struct join_fn : public join_base_fn { cudf::detail::input_offsetalator d_offsets; join_fn(column_device_view const d_strings, - string_view d_separator, - string_scalar_device_view d_narep) + column_device_view const d_separator, + column_device_view const d_narep) : join_base_fn{d_strings, d_separator, d_narep} { } @@ -104,8 +104,8 @@ struct join_fn : public join_base_fn { struct join_gather_fn : public join_base_fn { join_gather_fn(column_device_view const d_strings, - string_view d_separator, - string_scalar_device_view d_narep) + column_device_view const d_separator, + column_device_view const d_narep) : join_base_fn{d_strings, d_separator, d_narep} { } @@ -128,10 +128,12 @@ std::unique_ptr join_strings(strings_column_view const& input, { if (input.is_empty()) { return make_empty_column(type_id::STRING); } - CUDF_EXPECTS(separator.is_valid(stream), "Parameter separator must be a valid string_scalar"); + CUDF_EXPECTS(separator.is_valid(), "Parameter separator must be a valid string_scalar"); - string_view d_separator(separator.data(), separator.size()); - auto d_narep = get_scalar_device_view(const_cast(narep)); + auto const separator_view = separator.as_column_view(); + auto const narep_view = narep.as_column_view(); + auto const d_separator = column_device_view::create(separator_view.as_column_view(), stream); + auto const d_narep = column_device_view::create(narep_view.as_column_view(), stream); auto d_strings = column_device_view::create(input.parent(), stream); @@ -141,12 +143,12 @@ std::unique_ptr join_strings(strings_column_view const& input, ((input.chars_size(stream) / (input.size() - input.null_count())) <= AVG_CHAR_BYTES_THRESHOLD)) { return std::get<1>(make_strings_children( - join_fn{*d_strings, d_separator, d_narep}, input.size(), stream, mr)) + join_fn{*d_strings, *d_separator, *d_narep}, input.size(), stream, mr)) .release(); } // dynamically feeds index pairs to build the output auto indices = cudf::detail::make_counting_transform_iterator( - 0, join_gather_fn{*d_strings, d_separator, d_narep}); + 0, join_gather_fn{*d_strings, *d_separator, *d_narep}); auto joined_col = make_strings_column(indices, indices + (input.size() * 2), stream, mr); auto chars_data = joined_col->release().data; return std::move(*chars_data); @@ -164,7 +166,7 @@ std::unique_ptr join_strings(strings_column_view const& input, // build the null mask: only one output row so it is either all-valid or all-null auto const null_count = - static_cast(input.null_count() == input.size() && !narep.is_valid(stream)); + static_cast(input.null_count() == input.size() && !narep.is_valid()); auto null_mask = null_count ? cudf::detail::create_null_mask(1, cudf::mask_state::ALL_NULL, stream, mr) : rmm::device_buffer{0, stream, mr}; diff --git a/cpp/src/strings/combine/join_list_elements.cu b/cpp/src/strings/combine/join_list_elements.cu index 0a4c980ddbb1..5d73de0c7cb8 100644 --- a/cpp/src/strings/combine/join_list_elements.cu +++ b/cpp/src/strings/combine/join_list_elements.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include #include #include @@ -45,7 +45,7 @@ struct compute_size_and_concatenate_fn { column_device_view const lists_dv; size_type const* const list_offsets; column_device_view const strings_dv; - string_scalar_device_view const string_narep_dv; + column_device_view const string_narep_dv; separator_on_nulls const separate_nulls; output_if_empty_list const empty_list_policy; @@ -88,7 +88,7 @@ struct compute_size_and_concatenate_fn { bool null_element = strings_dv.is_null(str_idx); has_valid_element = has_valid_element || !null_element; - if (!d_chars && (null_element && !string_narep_dv.is_valid())) { + if (!d_chars && (null_element && !string_narep_dv.is_valid(0))) { size_bytes = 0; break; } @@ -99,8 +99,8 @@ struct compute_size_and_concatenate_fn { write_separator = false; } - auto const d_str = - null_element ? string_narep_dv.value() : strings_dv.element(str_idx); + auto const d_str = null_element ? string_narep_dv.element(0) + : strings_dv.element(str_idx); if (output_ptr) output_ptr = detail::copy_string(output_ptr, d_str); size_bytes += d_str.size_bytes(); @@ -120,7 +120,7 @@ struct compute_size_and_concatenate_fn { * separator is a string scalar. */ struct scalar_separator_fn { - string_scalar_device_view const d_separator; + column_device_view const d_separator; [[nodiscard]] __device__ bool is_null_list(column_device_view const& lists_dv, size_type const idx) const noexcept @@ -130,7 +130,7 @@ struct scalar_separator_fn { [[nodiscard]] __device__ string_view separator(size_type const) const noexcept { - return d_separator.value(); + return d_separator.element(0); } }; @@ -151,7 +151,7 @@ struct validities_fn { bool const valid_element = comp_fn.strings_dv.is_valid(str_idx); check_elements = check_elements || valid_element; // if an element is null and narep is invalid, the output row is null - if (!valid_element && !comp_fn.string_narep_dv.is_valid()) { return false; } + if (!valid_element && !comp_fn.string_narep_dv.is_valid(0)) { return false; } } // handle empty-list-as-null output policy setting valid_output = @@ -173,7 +173,7 @@ std::unique_ptr join_list_elements(lists_column_view const& lists_string { CUDF_EXPECTS(lists_strings_column.child().type().id() == type_id::STRING, "The input column must be a column of lists of strings"); - CUDF_EXPECTS(separator.is_valid(stream), "Parameter separator must be a valid string_scalar"); + CUDF_EXPECTS(separator.is_valid(), "Parameter separator must be a valid string_scalar"); auto const num_rows = lists_strings_column.size(); if (num_rows == 0) { return make_empty_column(type_id::STRING); } @@ -182,19 +182,22 @@ std::unique_ptr join_list_elements(lists_column_view const& lists_string // lists column, not `get_sliced_child()`. This is because calling to `offsets_begin()` on the // lists column returns a pointer to the offsets of the original lists column, which may not start // from `0`. - auto const strings_col = strings_column_view(lists_strings_column.child()); - auto const lists_dv_ptr = column_device_view::create(lists_strings_column.parent(), stream); - auto const strings_dv_ptr = column_device_view::create(strings_col.parent(), stream); - auto const sep_dv = get_scalar_device_view(const_cast(separator)); - auto const string_narep_dv = get_scalar_device_view(const_cast(narep)); - - auto const func = scalar_separator_fn{sep_dv}; + auto const strings_col = strings_column_view(lists_strings_column.child()); + auto const lists_dv_ptr = column_device_view::create(lists_strings_column.parent(), stream); + auto const strings_dv_ptr = column_device_view::create(strings_col.parent(), stream); + auto const separator_view = separator.as_column_view(); + auto const string_narep_view = narep.as_column_view(); + auto const sep_dv = column_device_view::create(separator_view.as_column_view(), stream); + auto const string_narep_dv = + column_device_view::create(string_narep_view.as_column_view(), stream); + + auto const func = scalar_separator_fn{*sep_dv}; auto const comp_fn = compute_size_and_concatenate_fn{func, *lists_dv_ptr, lists_strings_column.offsets_begin(), *strings_dv_ptr, - string_narep_dv, + *string_narep_dv, separate_nulls, empty_list_policy}; @@ -218,18 +221,18 @@ namespace { */ struct column_separators_fn { column_device_view const separators_dv; - string_scalar_device_view const sep_narep_dv; + column_device_view const sep_narep_dv; [[nodiscard]] __device__ bool is_null_list(column_device_view const& lists_dv, size_type const idx) const noexcept { - return lists_dv.is_null(idx) || (separators_dv.is_null(idx) && !sep_narep_dv.is_valid()); + return lists_dv.is_null(idx) || (separators_dv.is_null(idx) && !sep_narep_dv.is_valid(0)); } [[nodiscard]] __device__ string_view separator(size_type const idx) const noexcept { return separators_dv.is_valid(idx) ? separators_dv.element(idx) - : sep_narep_dv.value(); + : sep_narep_dv.element(0); } }; @@ -256,20 +259,23 @@ std::unique_ptr join_list_elements(lists_column_view const& lists_string // lists column, not `get_sliced_child()`. This is because calling to `offsets_begin()` on the // lists column returns a pointer to the offsets of the original lists column, which may not start // from `0`. - auto const strings_col = strings_column_view(lists_strings_column.child()); - auto const lists_dv_ptr = column_device_view::create(lists_strings_column.parent(), stream); - auto const strings_dv_ptr = column_device_view::create(strings_col.parent(), stream); - auto const string_narep_dv = get_scalar_device_view(const_cast(string_narep)); - auto const sep_dv_ptr = column_device_view::create(separators.parent(), stream); - auto const sep_narep_dv = get_scalar_device_view(const_cast(separator_narep)); - - auto const func = column_separators_fn{*sep_dv_ptr, sep_narep_dv}; + auto const strings_col = strings_column_view(lists_strings_column.child()); + auto const lists_dv_ptr = column_device_view::create(lists_strings_column.parent(), stream); + auto const strings_dv_ptr = column_device_view::create(strings_col.parent(), stream); + auto const string_narep_view = string_narep.as_column_view(); + auto const string_narep_dv = + column_device_view::create(string_narep_view.as_column_view(), stream); + auto const sep_dv_ptr = column_device_view::create(separators.parent(), stream); + auto const sep_narep_view = separator_narep.as_column_view(); + auto const sep_narep_dv = column_device_view::create(sep_narep_view.as_column_view(), stream); + + auto const func = column_separators_fn{*sep_dv_ptr, *sep_narep_dv}; auto const comp_fn = compute_size_and_concatenate_fn{func, *lists_dv_ptr, lists_strings_column.offsets_begin(), *strings_dv_ptr, - string_narep_dv, + *string_narep_dv, separate_nulls, empty_list_policy}; diff --git a/cpp/src/strings/filling/fill.cu b/cpp/src/strings/filling/fill.cu index badf27cf4d4d..94cdf40cfd98 100644 --- a/cpp/src/strings/filling/fill.cu +++ b/cpp/src/strings/filling/fill.cu @@ -4,6 +4,7 @@ */ #include +#include #include #include #include @@ -24,14 +25,14 @@ struct fill_fn { column_device_view const d_strings; size_type const begin; size_type const end; - string_scalar_device_view const d_value; + column_device_view const d_value; __device__ string_index_pair operator()(size_type idx) const { auto d_str = string_view(); if ((begin <= idx) && (idx < end)) { - if (!d_value.is_valid()) { return string_index_pair{nullptr, 0}; } - d_str = d_value.value(); + if (!d_value.is_valid(0)) { return string_index_pair{nullptr, 0}; } + d_str = d_value.element(0); } else { if (d_strings.is_null(idx)) { return string_index_pair{nullptr, 0}; } d_str = d_strings.element(idx); @@ -57,10 +58,11 @@ std::unique_ptr fill(strings_column_view const& input, CUDF_EXPECTS(begin <= end, "Parameters [begin,end) have invalid range values"); if (begin == end) { return std::make_unique(input.parent(), stream, mr); } - auto const d_strings = column_device_view::create(input.parent(), stream); - auto const d_value = cudf::get_scalar_device_view(const_cast(value)); + auto const d_strings = column_device_view::create(input.parent(), stream); + auto const value_view = value.as_column_view(); + auto const d_value = column_device_view::create(value_view.as_column_view(), stream); - auto fn = fill_fn{*d_strings, begin, end, d_value}; + auto fn = fill_fn{*d_strings, begin, end, *d_value}; rmm::device_uvector indices(strings_count, stream); thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), cuda::counting_iterator{0}, diff --git a/cpp/src/strings/search/contains_multiple.cu b/cpp/src/strings/search/contains_multiple.cu index 9f1d46b41049..d9d7b1364ba8 100644 --- a/cpp/src/strings/search/contains_multiple.cu +++ b/cpp/src/strings/search/contains_multiple.cu @@ -14,6 +14,7 @@ #include #include #include +#include #include #include #include diff --git a/cpp/src/strings/slice.cu b/cpp/src/strings/slice.cu index b154ac9b3374..58a2b68d3ab0 100644 --- a/cpp/src/strings/slice.cu +++ b/cpp/src/strings/slice.cu @@ -11,7 +11,6 @@ #include #include #include -#include #include #include #include diff --git a/cpp/src/text/normalize.cu b/cpp/src/text/normalize.cu index 110fd12b1306..00f0e7e2d42c 100644 --- a/cpp/src/text/normalize.cu +++ b/cpp/src/text/normalize.cu @@ -24,6 +24,7 @@ #include #include #include +#include #include #include diff --git a/cpp/tests/binaryop/binop-compiled-test.cpp b/cpp/tests/binaryop/binop-compiled-test.cpp index de60e0086b3b..2a06ae43dfc5 100644 --- a/cpp/tests/binaryop/binop-compiled-test.cpp +++ b/cpp/tests/binaryop/binop-compiled-test.cpp @@ -811,6 +811,36 @@ TEST_F(BinaryOperationCompiledTest_NullOpsString, NullMin_Vector_Vector) CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected, result->view()); } +TEST_F(BinaryOperationCompiledTest_NullOpsString, NullMinMax_ScalarOperands) +{ + auto const input = cudf::test::strings_column_wrapper( + {"aaa", "zzz", "unused", ""}, {true, true, false, true}); + auto const valid_scalar = cudf::string_scalar{"bbb"}; + auto const null_scalar = cudf::string_scalar{"unused", false}; + auto const string_type = cudf::data_type{cudf::type_id::STRING}; + + auto const expected_max = + cudf::test::strings_column_wrapper({"bbb", "zzz", "bbb", "bbb"}); + auto const expected_min = + cudf::test::strings_column_wrapper({"aaa", "bbb", "bbb", ""}); + + auto result = cudf::binary_operation( + input, valid_scalar, cudf::binary_operator::NULL_MAX, string_type); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_max, result->view()); + + result = cudf::binary_operation( + valid_scalar, input, cudf::binary_operator::NULL_MIN, string_type); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_min, result->view()); + + result = + cudf::binary_operation(null_scalar, input, cudf::binary_operator::NULL_MAX, string_type); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(input, result->view()); + + result = + cudf::binary_operation(input, null_scalar, cudf::binary_operator::NULL_MIN, string_type); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(input, result->view()); +} + TEST(BinaryOperationCompiledTest, LargeColumnNoOverflow) { cudf::size_type num_rows{1'799'989'091}; diff --git a/cpp/tests/dictionary/search_test.cpp b/cpp/tests/dictionary/search_test.cpp index fa1954eef061..527f275561f7 100644 --- a/cpp/tests/dictionary/search_test.cpp +++ b/cpp/tests/dictionary/search_test.cpp @@ -6,7 +6,11 @@ #include #include +#include +#include #include +#include +#include struct DictionarySearchTest : public cudf::test::BaseFixture {}; @@ -35,6 +39,20 @@ TEST_F(DictionarySearchTest, WithNulls) EXPECT_FALSE(result->is_valid()); } +TEST_F(DictionarySearchTest, FixedPointColumn) +{ + using decimal_type = numeric::decimal64; + auto const scale = numeric::scale_type{-2}; + auto const keys = cudf::test::fixed_point_column_wrapper{{100, 123, 250}, scale}; + auto const dictionary = cudf::dictionary::encode(keys); + auto const present = cudf::make_fixed_point_scalar(123, scale); + auto const missing = cudf::make_fixed_point_scalar(124, scale); + auto const dictionary_view = cudf::dictionary_column_view{dictionary->view()}; + + EXPECT_TRUE(cudf::dictionary::get_index(dictionary_view, *present)->is_valid()); + EXPECT_FALSE(cudf::dictionary::get_index(dictionary_view, *missing)->is_valid()); +} + TEST_F(DictionarySearchTest, EmptyColumn) { cudf::test::dictionary_column_wrapper dictionary{}; diff --git a/cpp/tests/iterator/scalar_iterator_test.cu b/cpp/tests/iterator/scalar_iterator_test.cu index 8a6936ab5952..6b4b221f5c92 100644 --- a/cpp/tests/iterator/scalar_iterator_test.cu +++ b/cpp/tests/iterator/scalar_iterator_test.cu @@ -35,10 +35,16 @@ TYPED_TEST(IteratorTest, scalar_iterator) [](auto v, auto b) { return cuda::std::pair{v, b}; }); // GPU test - auto it_dev = cudf::detail::make_scalar_iterator(*s); + auto const scalar_view = s->as_column_view(); + auto const d_scalar = + cudf::column_device_view::create(scalar_view.as_column_view(), cudf::get_default_stream()); + auto it_dev = cuda::make_permutation_iterator(d_scalar->begin(), + cuda::make_constant_iterator(0)); this->iterator_test_thrust(host_values, it_dev, host_values.size()); - auto it_pair_dev = cudf::detail::make_pair_iterator(*s); + auto it_pair_dev = + cuda::make_permutation_iterator(cudf::detail::make_pair_iterator(*d_scalar), + cuda::make_constant_iterator(0)); this->iterator_test_thrust(value_and_validity, it_pair_dev, host_values.size()); } @@ -49,11 +55,11 @@ TYPED_TEST(IteratorTest, null_scalar_iterator) cudf::test::UniformRandomGenerator(-128, 128).generate()); // data and valid arrays std::vector host_values(100, init); - std::vector host_bools(100, true); + std::vector host_bools(100, false); // create a scalar using ScalarType = cudf::scalar_type_t; - std::unique_ptr s(new ScalarType{init, true}); + std::unique_ptr s(new ScalarType{init, false}); // calculate the expected value by CPU. thrust::host_vector> value_and_validity(host_values.size()); @@ -64,6 +70,11 @@ TYPED_TEST(IteratorTest, null_scalar_iterator) [](auto v, auto b) { return cuda::std::pair{v, b}; }); // GPU test - auto it_pair_dev = cudf::detail::make_pair_iterator(*s); + auto const scalar_view = s->as_column_view(); + auto const d_scalar = + cudf::column_device_view::create(scalar_view.as_column_view(), cudf::get_default_stream()); + auto it_pair_dev = + cuda::make_permutation_iterator(cudf::detail::make_pair_iterator(*d_scalar), + cuda::make_constant_iterator(0)); this->iterator_test_thrust(value_and_validity, it_pair_dev, host_values.size()); } diff --git a/cpp/tests/iterator/value_iterator_test_strings.cu b/cpp/tests/iterator/value_iterator_test_strings.cu index e050229960b1..d5fcea40e793 100644 --- a/cpp/tests/iterator/value_iterator_test_strings.cu +++ b/cpp/tests/iterator/value_iterator_test_strings.cu @@ -129,9 +129,15 @@ TEST_F(StringIteratorTest, string_scalar_iterator) std::unique_ptr s(new ScalarType{zero, true}); // GPU test - auto it_dev = cudf::detail::make_scalar_iterator(*s); + auto const scalar_view = s->as_column_view(); + auto const d_scalar = + cudf::column_device_view::create(scalar_view.as_column_view(), cudf::get_default_stream()); + auto it_dev = cuda::make_permutation_iterator(d_scalar->begin(), + cuda::make_constant_iterator(0)); this->iterator_test_thrust(all_array, it_dev, host_values.size()); - auto it_pair_dev = cudf::detail::make_pair_iterator(*s); + auto it_pair_dev = + cuda::make_permutation_iterator(cudf::detail::make_pair_iterator(*d_scalar), + cuda::make_constant_iterator(0)); this->iterator_test_thrust(value_and_validity, it_pair_dev, host_values.size()); } diff --git a/cpp/tests/jit/row_ir.cpp b/cpp/tests/jit/row_ir.cpp index e6ee4c7b12f6..cc30f6b6b471 100644 --- a/cpp/tests/jit/row_ir.cpp +++ b/cpp/tests/jit/row_ir.cpp @@ -356,8 +356,7 @@ TEST_F(RowIRCudaCodeGenTest, AstConversionBasic) cudf::get_default_stream(), cudf::get_current_device_resource_ref()); - ASSERT_EQ(transform_args.scalar_columns.size(), 1); - ASSERT_EQ(transform_args.scalar_columns[0]->view().size(), 1); + ASSERT_EQ(transform_args.scalar_columns.size(), 0); EXPECT_EQ(transform_args.source_type, cudf::udf_source_type::CUDA); EXPECT_EQ(transform_args.is_null_aware, cudf::null_aware::NO); EXPECT_EQ(transform_args.outputs.size(), 1); @@ -370,6 +369,8 @@ TEST_F(RowIRCudaCodeGenTest, AstConversionBasic) EXPECT_EQ(std::get(transform_args.inputs[0]).type(), cudf::data_type{cudf::type_id::INT32}); EXPECT_EQ(std::get(transform_args.inputs[0]).null_count(), 0); + EXPECT_EQ(std::get(transform_args.inputs[0]).data(), + forty_two.data()); /// The input column should be the second column in the transform args ASSERT_TRUE(std::holds_alternative(transform_args.inputs[1])); diff --git a/cpp/tests/scalar/scalar_device_view_test.cu b/cpp/tests/scalar/scalar_device_view_test.cu index b20f5812c8de..5f00c9cb48b2 100644 --- a/cpp/tests/scalar/scalar_device_view_test.cu +++ b/cpp/tests/scalar/scalar_device_view_test.cu @@ -31,7 +31,6 @@ template CUDF_KERNEL void test_set_value(ScalarDeviceViewType s, ScalarDeviceViewType s1) { s1.set_value(s.value()); - s1.set_valid(true); } template @@ -84,26 +83,6 @@ TYPED_TEST(TypedScalarDeviceViewTest, ConstructNull) EXPECT_FALSE(result.value(cudf::get_default_stream())); } -template -CUDF_KERNEL void test_setnull(ScalarDeviceViewType s) -{ - s.set_valid(false); -} - -TYPED_TEST(TypedScalarDeviceViewTest, SetNull) -{ - TypeParam value = cudf::test::make_type_param_scalar(5); - cudf::scalar_type_t s{value}; - auto scalar_device_view = cudf::get_scalar_device_view(s); - s.set_valid_async(true); - EXPECT_TRUE(s.is_valid()); - - test_setnull<<<1, 1, 0, cudf::get_default_stream().value()>>>(scalar_device_view); - CUDF_CHECK_CUDA(0); - - EXPECT_FALSE(s.is_valid()); -} - struct StringScalarDeviceViewTest : public cudf::test::BaseFixture {}; CUDF_KERNEL void test_string_value(cudf::string_scalar_device_view s, diff --git a/cpp/tests/scalar/scalar_test.cpp b/cpp/tests/scalar/scalar_test.cpp index c7b4f2f1de02..ccf8e8aacfc0 100644 --- a/cpp/tests/scalar/scalar_test.cpp +++ b/cpp/tests/scalar/scalar_test.cpp @@ -9,7 +9,10 @@ #include #include +#include +#include #include +#include #include #include @@ -62,7 +65,7 @@ class lifetime_test_scalar : public cudf::numeric_scalar { void set_data_async(int32_t const& value, cuda::stream_ref stream) { - this->_data.set_value_async(value, stream); + this->set_value(value, stream); } }; @@ -226,6 +229,13 @@ TEST_F(StringScalarTest, MoveConstructor) EXPECT_EQ(data_ptr, s2.data()); } +TEST_F(StringScalarTest, OverflowCheckedBeforeAllocation) +{ + char source{}; + auto const oversized = static_cast(std::numeric_limits::max()) + 1; + EXPECT_THROW(cudf::string_scalar(std::string_view{&source, oversized}), std::overflow_error); +} + struct ListScalarTest : public cudf::test::BaseFixture {}; TEST_F(ListScalarTest, DefaultValidityNonNested) @@ -246,6 +256,76 @@ TEST_F(ListScalarTest, DefaultValidityNested) CUDF_TEST_EXPECT_COLUMNS_EQUAL(data, s.view()); } +TEST_F(ScalarTest, OneRowColumnLayoutBaseline) +{ + auto const numeric = cudf::numeric_scalar{42}; + auto const null_numeric = cudf::numeric_scalar{42, false}; + auto const decimal = cudf::fixed_point_scalar{1234, numeric::scale_type{-2}}; + auto const decimal128 = + cudf::fixed_point_scalar{__int128_t{5678}, numeric::scale_type{-7}}; + auto const raw_int128 = cudf::numeric_scalar<__int128_t>{__int128_t{9012}}; + auto const string = cudf::string_scalar{"scalar"}; + auto const null_string = cudf::string_scalar{"null scalar", false}; + auto const list_elements = cudf::test::fixed_width_column_wrapper{1, 2, 3}; + auto const list = cudf::list_scalar{list_elements}; + auto const null_list = cudf::list_scalar{list_elements, false}; + auto const nested_elements = cudf::test::lists_column_wrapper{{1, 2}, {}, {3}}; + auto const nested_list = cudf::list_scalar{nested_elements}; + + auto const numeric_column = numeric.as_column_view(); + auto const null_numeric_column = null_numeric.as_column_view(); + auto const decimal_column = decimal.as_column_view(); + auto const decimal128_column = decimal128.as_column_view(); + auto const raw_int128_column = raw_int128.as_column_view(); + auto const string_column = string.as_column_view(); + auto const null_string_column = null_string.as_column_view(); + auto const list_column = list.as_column_view(); + auto const null_list_column = null_list.as_column_view(); + auto const nested_list_column = nested_list.as_column_view(); + + auto const expected_numeric = cudf::test::fixed_width_column_wrapper{42}; + auto const expected_decimal = + cudf::test::fixed_point_column_wrapper({1234}, numeric::scale_type{-2}); + auto const expected_decimal128 = + cudf::test::fixed_point_column_wrapper<__int128_t>({5678}, numeric::scale_type{-7}); + auto const expected_raw_int128 = + cudf::test::fixed_point_column_wrapper<__int128_t>({9012}, numeric::scale_type{0}); + auto const expected_string = cudf::test::strings_column_wrapper{"scalar"}; + auto const expected_list = cudf::test::lists_column_wrapper{{1, 2, 3}}; + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_numeric, numeric_column.as_column_view()); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_decimal, decimal_column.as_column_view()); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_decimal128, decimal128_column.as_column_view()); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_raw_int128, raw_int128_column.as_column_view()); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_string, string_column.as_column_view()); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_list, list_column.as_column_view()); + + EXPECT_EQ(1, null_numeric_column.null_count()); + EXPECT_TRUE(null_numeric_column.nullable()); + EXPECT_EQ(cudf::data_type(cudf::type_id::DECIMAL128, 0), raw_int128_column.type()); + EXPECT_EQ(1, null_string_column.null_count()); + EXPECT_EQ(1, null_list_column.null_count()); + EXPECT_EQ(cudf::type_id::STRING, null_string_column.type().id()); + EXPECT_EQ(cudf::type_id::LIST, null_list_column.type().id()); + EXPECT_EQ(2, null_string_column.as_column_view().child(0).size()); + EXPECT_EQ(3, null_list_column.as_column_view().child(1).size()); + EXPECT_EQ(cudf::type_id::LIST, nested_list_column.as_column_view().child(1).type().id()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(nested_elements, nested_list_column.as_column_view().child(1)); + EXPECT_EQ(numeric.data(), numeric_column.data()); + EXPECT_EQ(string.data(), string_column.data()); + EXPECT_EQ(list.view().data(), list_column.as_column_view().child(1).data()); +} + +TEST_F(ScalarTest, MutableOneRowColumnLayout) +{ + auto numeric = cudf::numeric_scalar{42}; + auto mutable_column = numeric.as_mutable_column_view(); + + EXPECT_EQ(1, mutable_column.size()); + EXPECT_EQ(numeric.data(), mutable_column.data()); + EXPECT_EQ(numeric.validity_data(), mutable_column.null_mask()); +} + TEST_F(ListScalarTest, MoveColumnConstructor) { auto data = cudf::test::fixed_width_column_wrapper{1, 2, 3}; diff --git a/cpp/tests/search/search_test.cpp b/cpp/tests/search/search_test.cpp index f7bfefad5e42..8d55d8c3f5a0 100644 --- a/cpp/tests/search/search_test.cpp +++ b/cpp/tests/search/search_test.cpp @@ -10,6 +10,7 @@ #include #include +#include #include #include @@ -581,6 +582,18 @@ TEST_F(SearchTest, contains_false) ASSERT_EQ(result, expect); } +TEST_F(SearchTest, contains_fixed_point) +{ + using decimal_type = numeric::decimal64; + auto const scale = numeric::scale_type{-2}; + auto const column = cudf::test::fixed_point_column_wrapper{{100, 123, 250}, scale}; + auto const present = cudf::make_fixed_point_scalar(123, scale); + auto const missing = cudf::make_fixed_point_scalar(124, scale); + + EXPECT_TRUE(cudf::contains(column, *present)); + EXPECT_FALSE(cudf::contains(column, *missing)); +} + TEST_F(SearchTest, contains_empty_value) { using element_type = int64_t;