diff --git a/cpp/include/cudf/detail/unary.hpp b/cpp/include/cudf/detail/unary.hpp index 285a798a4787..e55ccfb16aea 100644 --- a/cpp/include/cudf/detail/unary.hpp +++ b/cpp/include/cudf/detail/unary.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2018-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -8,50 +8,12 @@ #include #include #include -#include #include #include -#include -#include - -namespace CUDF_EXPORT cudf { +namespace cudf { namespace detail { -/** - * @brief Creates a column of `type_id::BOOL8` elements by applying a predicate to every element - * between - * [`begin, `end`) `true` indicates the value is satisfies the predicate and `false` indicates it - * doesn't. - * - * @tparam InputIterator Iterator type for `begin` and `end` - * @tparam Predicate A predicator type which will be evaluated - * @param begin Beginning of the sequence of elements - * @param end End of the sequence of elements - * @param p Predicate to be applied to each element in `[begin,end)` - * @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 - * - * @returns A column of type `type_id::BOOL8,` with `true` representing predicate is satisfied. - */ - -template -std::unique_ptr true_if(InputIterator begin, - InputIterator end, - size_type size, - Predicate p, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) -{ - auto output = - make_numeric_column(data_type(type_id::BOOL8), size, mask_state::UNALLOCATED, stream, mr); - auto output_mutable_view = output->mutable_view(); - auto output_data = output_mutable_view.data(); - - thrust::transform(rmm::exec_policy_nosync(stream), begin, end, output_data, p); - - return output; -} /** * @copydoc cudf::unary_operation @@ -91,4 +53,4 @@ std::unique_ptr is_not_nan(cudf::column_view const& input, rmm::device_async_resource_ref mr); } // namespace detail -} // namespace CUDF_EXPORT cudf +} // namespace cudf diff --git a/cpp/src/unary/math_ops.cu b/cpp/src/unary/math_ops.cu index 742735e848d5..8aae1c01d85b 100644 --- a/cpp/src/unary/math_ops.cu +++ b/cpp/src/unary/math_ops.cu @@ -15,6 +15,7 @@ #include #include +#include #include #include diff --git a/cpp/src/unary/nan_ops.cu b/cpp/src/unary/nan_ops.cu index c5be3d038342..5cfaa0e170d7 100644 --- a/cpp/src/unary/nan_ops.cu +++ b/cpp/src/unary/nan_ops.cu @@ -1,8 +1,10 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ +#include "true_if.cuh" + #include #include #include diff --git a/cpp/src/unary/null_ops.cu b/cpp/src/unary/null_ops.cu index d38427cc0d28..cff98f234146 100644 --- a/cpp/src/unary/null_ops.cu +++ b/cpp/src/unary/null_ops.cu @@ -1,8 +1,10 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ +#include "true_if.cuh" + #include #include #include diff --git a/cpp/src/unary/true_if.cuh b/cpp/src/unary/true_if.cuh new file mode 100644 index 000000000000..90bd4290e3a0 --- /dev/null +++ b/cpp/src/unary/true_if.cuh @@ -0,0 +1,55 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include + +#include +#include + +#include + +namespace cudf { +namespace detail { +/** + * @brief Creates a column of `type_id::BOOL8` elements by applying a predicate to every element + * between + * [`begin, `end`) `true` indicates the value is satisfies the predicate and `false` indicates it + * doesn't. + * + * @tparam InputIterator Iterator type for `begin` and `end` + * @tparam Predicate A predicator type which will be evaluated + * @param begin Beginning of the sequence of elements + * @param end End of the sequence of elements + * @param size Size of the output column + * @param p Predicate to be applied to each element in `[begin,end)` + * @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 + * + * @returns A column of type `type_id::BOOL8,` with `true` representing predicate is satisfied. + */ + +template +std::unique_ptr true_if(InputIterator begin, + InputIterator end, + size_type size, + Predicate p, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto output = + make_numeric_column(data_type(type_id::BOOL8), size, mask_state::UNALLOCATED, stream, mr); + auto output_mutable_view = output->mutable_view(); + auto output_data = output_mutable_view.data(); + + thrust::transform(rmm::exec_policy_nosync(stream), begin, end, output_data, p); + + return output; +} + +} // namespace detail +} // namespace cudf