diff --git a/cpp/include/cudf/detail/null_mask.cuh b/cpp/include/cudf/detail/null_mask.cuh index 456de20cfe25..6159bf0c129a 100644 --- a/cpp/include/cudf/detail/null_mask.cuh +++ b/cpp/include/cudf/detail/null_mask.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -224,6 +224,27 @@ CUDF_KERNEL void segmented_offset_bitmask_binop(Binop op, if (lane == 0) { null_counts[segment_id] = warp_count; } } +// Forward declarations; defined later in this header but called from the templates below. +template +size_type inplace_bitmask_binop(Binop op, + device_span dest_mask, + host_span masks, + host_span masks_begin_bits, + size_type mask_size_bits, + rmm::cuda_stream_view stream); + +template +rmm::device_uvector inplace_segmented_bitmask_binop( + Binop op, + device_span dest_masks, + size_type dest_mask_size, + host_span masks, + host_span masks_begin_bits, + size_type mask_size_bits, + host_span segment_offsets, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); + /** * @copydoc bitmask_binop(Binop op, host_span, host_span * const, size_type, rmm::device_async_resource_ref) diff --git a/cpp/include/cudf/utilities/span.hpp b/cpp/include/cudf/utilities/span.hpp index aace1adff144..457b448aff26 100644 --- a/cpp/include/cudf/utilities/span.hpp +++ b/cpp/include/cudf/utilities/span.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -19,7 +19,6 @@ #include #include -#include #include #include #include @@ -33,131 +32,7 @@ namespace CUDF_EXPORT cudf { */ /// A constant used to differentiate std::span of static and dynamic extent -constexpr std::size_t dynamic_extent = std::numeric_limits::max(); - -/** @} */ // end of group -namespace detail { - -/** - * @brief C++20 std::span with reduced feature set. - * - */ -template -class span_base { - static_assert(Extent == dynamic_extent, "Only dynamic extent is supported"); - - public: - using element_type = T; ///< The type of the elements in the span - using value_type = std::remove_cv; ///< Stored value type - using size_type = std::size_t; ///< The type used for the size of the span - using difference_type = std::ptrdiff_t; ///< std::ptrdiff_t - using pointer = T*; ///< The type of the pointer returned by data() - using iterator = T*; ///< The type of the iterator returned by begin() - using const_pointer = T const*; ///< The type of the pointer returned by data() const - using reference = T&; ///< The type of the reference returned by operator[](size_type) - using const_reference = - T const&; ///< The type of the reference returned by operator[](size_type) const - - static constexpr std::size_t extent = Extent; ///< The extent of the span - - CUDF_HOST_DEVICE constexpr span_base() noexcept {} - /** - * @brief Constructs a span from a pointer and a size. - * - * @param data Pointer to the first element in the span. - * @param size The number of elements in the span. - */ - CUDF_HOST_DEVICE constexpr span_base(pointer data, size_type size) : _data(data), _size(size) {} - // constexpr span_base(pointer begin, pointer end) : _data(begin), _size(end - begin) {} - constexpr span_base(span_base const&) noexcept = default; ///< Copy constructor - /** - * @brief Copy assignment operator. - * - * @return Reference to this span. - */ - constexpr span_base& operator=(span_base const&) noexcept = default; - - /** - * @brief Returns an iterator to the first element of the span. - * - * If the span is empty, the returned iterator will be equal to end(). - * - * @return An iterator to the first element of the span - */ - [[nodiscard]] CUDF_HOST_DEVICE constexpr iterator begin() const noexcept { return _data; } - /** - * @brief Returns an iterator to the element following the last element of the span. - * - * This element acts as a placeholder; attempting to access it results in undefined behavior. - * - * @return An iterator to the element following the last element of the span - */ - [[nodiscard]] CUDF_HOST_DEVICE constexpr iterator end() const noexcept { return _data + _size; } - /** - * @brief Returns a pointer to the beginning of the sequence. - * - * @return A pointer to the first element of the span - */ - [[nodiscard]] CUDF_HOST_DEVICE constexpr pointer data() const noexcept { return _data; } - - /** - * @brief Returns the number of elements in the span. - * - * @return The number of elements in the span - */ - [[nodiscard]] CUDF_HOST_DEVICE constexpr size_type size() const noexcept { return _size; } - /** - * @brief Returns the size of the sequence in bytes. - * - * @return The size of the sequence in bytes - */ - [[nodiscard]] CUDF_HOST_DEVICE constexpr size_type size_bytes() const noexcept - { - return sizeof(T) * _size; - } - - /** - * @brief Checks if the span is empty. - * - * @return True if the span is empty, false otherwise - */ - [[nodiscard]] CUDF_HOST_DEVICE constexpr bool empty() const noexcept { return _size == 0; } - - /** - * @brief Obtains a subspan consisting of the first N elements of the sequence - * - * @param count Number of elements from the beginning of this span to put in the subspan. - * @return A subspan of the first N elements of the sequence - */ - [[nodiscard]] constexpr Derived first(size_type count) const noexcept - { - return Derived(_data, count); - } - - /** - * @brief Obtains a subspan consisting of the last N elements of the sequence - * - * @param count Number of elements from the end of this span to put in the subspan - * @return A subspan of the last N elements of the sequence - */ - [[nodiscard]] constexpr Derived last(size_type count) const noexcept - { - return Derived(_data + _size - count, count); - } - - protected: - pointer _data{nullptr}; ///< Pointer to the first element in the span - size_type _size{0}; ///< The number of elements in the span -}; - -} // namespace detail - -/** - * @addtogroup utility_span - * @{ - * @file - * @brief APIs for spans - */ +constexpr std::size_t dynamic_extent = cuda::std::dynamic_extent; // ===== host_span ================================================================================= @@ -177,18 +52,44 @@ struct is_host_span_supported_container< // std::basic_string, Alloc>> : std::true_type {}; /** - * @brief C++20 std::span with reduced feature set. + * @brief Host span, a non-owning view over a contiguous sequence of host-accessible elements. * + * Backed by `cuda::std::span`, with additional support for constructing from cudf-supported host + * containers and for tracking whether the underlying memory is device accessible (e.g. pinned + * memory), which enables copy engine optimizations. */ template -struct host_span : public cudf::detail::span_base> { - using base = cudf::detail::span_base>; ///< Base type - using base::base; +struct host_span { + private: + using span_type = cuda::std::span; ///< The underlying span type + + public: + using element_type = typename span_type::element_type; ///< Element type + using value_type = typename span_type::value_type; ///< Stored value type + using size_type = typename span_type::size_type; ///< Size type + using difference_type = typename span_type::difference_type; ///< std::ptrdiff_t + using pointer = typename span_type::pointer; ///< Pointer returned by data() + using const_pointer = typename span_type::const_pointer; ///< Pointer returned by data() const + using reference = typename span_type::reference; ///< Reference returned by operator[] + using const_reference = typename span_type::const_reference; ///< Const reference to an element + using iterator = pointer; ///< The type of the iterator returned by begin() - constexpr host_span() noexcept : base() {} // required to compile on centos + static constexpr std::size_t extent = span_type::extent; ///< The extent of the span + + constexpr host_span() noexcept {} // required to compile on centos + + /** + * @brief Constructs a span from a pointer and a size. + * + * @note This needs to be host-device, as it's used by a host-device function in base_2dspan + * + * @param data Pointer to the first element in the span + * @param size The number of elements in the span + */ + CUDF_HOST_DEVICE constexpr host_span(T* data, std::size_t size) : _span{data, size} {} /** - * @brief Constructor from pointer and size + * @brief Constructor from pointer, size and device-accessibility flag * * @note This needs to be host-device, as it's used by a host-device function in base_2dspan * @@ -197,7 +98,7 @@ struct host_span : public cudf::detail::span_base().data()))> (*)[], T (*)[]>>* = nullptr> // NOLINT - constexpr host_span(C& in) : base(thrust::raw_pointer_cast(in.data()), in.size()) + constexpr host_span(C& in) : _span{thrust::raw_pointer_cast(in.data()), in.size()} { } @@ -223,7 +124,7 @@ struct host_span : public cudf::detail::span_base().data()))> (*)[], T (*)[]>>* = nullptr> // NOLINT - constexpr host_span(C const& in) : base(thrust::raw_pointer_cast(in.data()), in.size()) + constexpr host_span(C const& in) : _span{thrust::raw_pointer_cast(in.data()), in.size()} { } @@ -235,9 +136,10 @@ struct host_span : public cudf::detail::span_base, // NOLINT void>* = nullptr> constexpr host_span(host_span const& other) noexcept - : base(other.data(), other.size()), _is_device_accessible{other.is_device_accessible()} + : _span{other.data(), other.size()}, _is_device_accessible{other.is_device_accessible()} { } + // not noexcept due to undefined behavior when idx < 0 || idx >= size /** * @brief Returns a reference to the idx-th element of the sequence. @@ -248,10 +150,10 @@ struct host_span : public cudf::detail::span_base= sizeof(size_t), "index type must not be smaller than size_t"); - return this->_data[idx]; + return _span[idx]; } // not noexcept due to undefined behavior when size = 0 @@ -262,18 +164,86 @@ struct host_span : public cudf::detail::span_base_data[0]; } + [[nodiscard]] constexpr reference front() const { return _span.front(); } // not noexcept due to undefined behavior when size = 0 /** * @brief Returns a reference to the last element in the span. * - * Calling last on an empty span results in undefined behavior. + * Calling back on an empty span results in undefined behavior. * * @return Reference to the last element in the span */ - [[nodiscard]] constexpr typename base::reference back() const + [[nodiscard]] constexpr reference back() const { return _span.back(); } + + /** + * @brief Returns an iterator to the first element of the span. + * + * If the span is empty, the returned iterator will be equal to end(). + * + * @return An iterator to the first element of the span + */ + [[nodiscard]] CUDF_HOST_DEVICE constexpr iterator begin() const noexcept { return _span.data(); } + /** + * @brief Returns an iterator to the element following the last element of the span. + * + * This element acts as a placeholder; attempting to access it results in undefined behavior. + * + * @return An iterator to the element following the last element of the span + */ + [[nodiscard]] CUDF_HOST_DEVICE constexpr iterator end() const noexcept + { + return _span.data() + _span.size(); + } + /** + * @brief Returns a pointer to the beginning of the sequence. + * + * @return A pointer to the first element of the span + */ + [[nodiscard]] CUDF_HOST_DEVICE constexpr pointer data() const noexcept { return _span.data(); } + + /** + * @brief Returns the number of elements in the span. + * + * @return The number of elements in the span + */ + [[nodiscard]] CUDF_HOST_DEVICE constexpr size_type size() const noexcept { return _span.size(); } + /** + * @brief Returns the size of the sequence in bytes. + * + * @return The size of the sequence in bytes + */ + [[nodiscard]] CUDF_HOST_DEVICE constexpr size_type size_bytes() const noexcept + { + return _span.size_bytes(); + } + + /** + * @brief Checks if the span is empty. + * + * @return True if the span is empty, false otherwise + */ + [[nodiscard]] CUDF_HOST_DEVICE constexpr bool empty() const noexcept { return _span.empty(); } + + /** + * @brief Obtains a subspan consisting of the first count elements of the sequence + * + * @param count Number of elements from the beginning of this span to put in the subspan. + * @return A subspan of the first count elements of the sequence + */ + [[nodiscard]] constexpr host_span first(size_type count) const noexcept + { + return host_span{_span.data(), count, _is_device_accessible}; + } + + /** + * @brief Obtains a subspan consisting of the last count elements of the sequence + * + * @param count Number of elements from the end of this span to put in the subspan + * @return A subspan of the last count elements of the sequence + */ + [[nodiscard]] constexpr host_span last(size_type count) const noexcept { - return this->_data[this->_size - 1]; + return host_span{_span.data() + _span.size() - count, count, _is_device_accessible}; } /** @@ -290,10 +260,10 @@ struct host_span : public cudf::detail::span_basedata() + offset, count, _is_device_accessible}; + return host_span{_span.data() + offset, count, _is_device_accessible}; } /** @@ -303,10 +273,13 @@ struct host_span : public cudf::detail::span_base() const noexcept { - return std::span(this->data(), this->size()); + return std::span(_span.data(), _span.size()); } private: + // TODO: could be std::span once base_2dspan moves to cuda::std::mdspan and host_span no longer + // needs to be device-usable. + span_type _span; bool _is_device_accessible{false}; }; diff --git a/docs/cudf/source/conf.py b/docs/cudf/source/conf.py index c73e4b43c839..191b7d8df59f 100644 --- a/docs/cudf/source/conf.py +++ b/docs/cudf/source/conf.py @@ -457,6 +457,8 @@ def _generate_namespaces(namespaces): "orc::column_statistics", # Span subclasses access base class members "base::", + # host_span defines member typedefs via its underlying cuda::std::span alias + "span_type", } _domain_objects = None