Move true_if device function from detail/.hpp to src/.cuh header

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 <cudf/column/column_factories.hpp>
 #include <cudf/unary.hpp>
 #include <cudf/utilities/default_stream.hpp>
-#include <cudf/utilities/export.hpp>
 #include <cudf/utilities/memory_resource.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/exec_policy.hpp>
 
-#include <thrust/transform.h>
-
-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 <typename InputIterator, typename Predicate>
-std::unique_ptr<column> 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<bool>();
-
-  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<column> is_not_nan(cudf::column_view const& input,
                                    rmm::device_async_resource_ref mr);
 
 }  // namespace detail
-}  // namespace CUDF_EXPORT cudf
+}  // namespace cudf

cpp/src/unary/math_ops.cu

@@ -15,6 +15,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <cuda/std/bit>
 #include <cuda/std/cmath>

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 <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/iterator.cuh>

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 <cudf/column/column_device_view.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/unary.hpp>

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 <cudf/column/column_factories.hpp>
+#include <cudf/utilities/memory_resource.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <thrust/transform.h>
+
+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 <typename InputIterator, typename Predicate>
+std::unique_ptr<column> 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<bool>();
+
+  thrust::transform(rmm::exec_policy_nosync(stream), begin, end, output_data, p);
+
+  return output;
+}
+
+}  // namespace detail
+}  // namespace cudf