Move compute_column API out of ast namespace

conda/recipes/libcudf/meta.yaml

@@ -51,7 +51,6 @@ test:
     - test -f $PREFIX/lib/libcudf.so
     - test -f $PREFIX/lib/libcudftestutil.a
     - test -f $PREFIX/include/cudf/aggregation.hpp
-    - test -f $PREFIX/include/cudf/ast/transform.hpp
     - test -f $PREFIX/include/cudf/ast/detail/expression_parser.hpp
     - test -f $PREFIX/include/cudf/ast/detail/operators.hpp
     - test -f $PREFIX/include/cudf/ast/nodes.hpp

cpp/CMakeLists.txt

@@ -152,7 +152,6 @@ add_library(cudf
     src/aggregation/aggregation.cu
     src/aggregation/result_cache.cpp
     src/ast/expression_parser.cpp
-    src/ast/transform.cu
     src/binaryop/binaryop.cpp
     src/binaryop/compiled/binary_ops.cu
     src/binaryop/compiled/Add.cu
@@ -436,6 +435,7 @@ add_library(cudf
     src/text/subword/wordpiece_tokenizer.cu
     src/text/tokenize.cu
     src/transform/bools_to_mask.cu
+    src/transform/compute_column.cu
     src/transform/encode.cu
     src/transform/mask_to_bools.cu
     src/transform/nans_to_nulls.cu

cpp/benchmarks/ast/transform_benchmark.cpp

@@ -14,10 +14,10 @@
  * limitations under the License.
  */
 
-#include <cudf/ast/transform.hpp>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/table/table.hpp>
 #include <cudf/table/table_view.hpp>
+#include <cudf/transform.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/error.hpp>
 
@@ -119,7 +119,7 @@ static void BM_ast_transform(benchmark::State& state)
   // Execute benchmark
   for (auto _ : state) {
     cuda_event_timer raii(state, true);  // flush_l2_cache = true, stream = 0
-    cudf::ast::compute_column(table, expression_tree_root);
+    cudf::compute_column(table, expression_tree_root);
   }
 
   // Use the number of bytes read from global memory

cpp/include/cudf/ast/detail/expression_evaluator.cuh

@@ -191,13 +191,13 @@ struct mutable_column_expression_result
 };
 
 /**
- * @brief Despite to a binary operator based on a single data type.
+ * @brief Dispatch to a binary operator based on a single data type.
  *
  * This functor is a dispatcher for binary operations that assumes that both
- * operands to a binary operation are of the same type. This assumption is
- * encoded in the one non-deducible template parameter LHS, the type of the
- * left-hand operand, which is then used as the template parameter for both the
- * left and right operands to the binary operator f.
+ * operands are of the same type. This assumption is encoded in the
+ * non-deducible template parameter LHS, the type of the left-hand operand,
+ * which is then used as the template parameter for both the left and right
+ * operands to the binary operator f.
  */
 struct single_dispatch_binary_operator {
   /**

cpp/include/cudf/ast/detail/expression_parser.hpp

@@ -30,7 +30,7 @@ namespace ast {
 namespace detail {
 
 /**
- * @brief Enum defining data reference types used by a node.
+ * @brief Node data reference types.
  *
  * This enum is device-specific. For instance, intermediate data references are generated by the
  * linearization process but cannot be explicitly created by the user.

cpp/include/cudf/detail/transform.hpp

@@ -16,6 +16,7 @@
 
 #pragma once
 
+#include <cudf/ast/nodes.hpp>
 #include <cudf/transform.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
@@ -35,6 +36,17 @@ std::unique_ptr<column> transform(
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+/**
+ * @copydoc cudf::compute_column
+ *
+ * @param stream CUDA stream used for device memory operations and kernel launches.
+ */
+std::unique_ptr<column> compute_column(
+  table_view const table,
+  ast::expression const& expr,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 /**
  * @copydoc cudf::nans_to_nulls
  *

cpp/include/cudf/transform.hpp

@@ -16,6 +16,7 @@
 
 #pragma once
 
+#include <cudf/ast/nodes.hpp>
 #include <cudf/types.hpp>
 
 #include <memory>
@@ -74,6 +75,22 @@ std::pair<std::unique_ptr<rmm::device_buffer>, size_type> nans_to_nulls(
   column_view const& input,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+/**
+ * @brief Compute a new column by evaluating an expression tree on a table.
+ *
+ * This evaluates an expression over a table to produce a new column. Also called an n-ary
+ * transform.
+ *
+ * @param table The table used for expression evaluation.
+ * @param expr The root of the expression tree.
+ * @param mr Device memory resource.
+ * @return std::unique_ptr<column> Output column.
+ */
+std::unique_ptr<column> compute_column(
+  table_view const table,
+  ast::expression const& expr,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 /**
  * @brief Creates a bitmask from a column of boolean elements.
  *

cpp/src/join/conditional_join.cuh

@@ -19,7 +19,7 @@
 #include <join/join_common_utils.cuh>
 #include <join/join_common_utils.hpp>
 
-#include <cudf/ast/detail/transform.cuh>
+#include <cudf/ast/detail/expression_parser.hpp>
 #include <cudf/ast/nodes.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/table/table.hpp>

cpp/src/ast/transform.cu → cpp/src/transform/compute_column.cu

@@ -18,14 +18,15 @@
 #include <cudf/ast/detail/expression_parser.hpp>
 #include <cudf/ast/nodes.hpp>
 #include <cudf/ast/operators.hpp>
-#include <cudf/ast/transform.hpp>
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/transform.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/scalar/scalar_device_view.cuh>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/table/table_view.hpp>
+#include <cudf/transform.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/error.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
@@ -34,7 +35,6 @@
 #include <rmm/mr/device/device_memory_resource.hpp>
 
 namespace cudf {
-namespace ast {
 namespace detail {
 
 /**
@@ -63,8 +63,8 @@ __launch_bounds__(max_block_size) __global__
   // workaround is to declare an arbitrary (here char) array type then cast it
   // after the fact to the appropriate type.
   extern __shared__ char raw_intermediate_storage[];
-  IntermediateDataType<has_nulls>* intermediate_storage =
-    reinterpret_cast<IntermediateDataType<has_nulls>*>(raw_intermediate_storage);
+  ast::detail::IntermediateDataType<has_nulls>* intermediate_storage =
+    reinterpret_cast<ast::detail::IntermediateDataType<has_nulls>*>(raw_intermediate_storage);
 
   auto thread_intermediate_storage =
     &intermediate_storage[threadIdx.x * device_expression_data.num_intermediates];
@@ -74,13 +74,13 @@ __launch_bounds__(max_block_size) __global__
     table, device_expression_data, thread_intermediate_storage);
 
   for (cudf::size_type row_index = start_idx; row_index < table.num_rows(); row_index += stride) {
-    auto output_dest = mutable_column_expression_result<has_nulls>(output_column);
+    auto output_dest = ast::detail::mutable_column_expression_result<has_nulls>(output_column);
     evaluator.evaluate(output_dest, row_index);
   }
 }
 
 std::unique_ptr<column> compute_column(table_view const table,
-                                       expression const& expr,
+                                       ast::expression const& expr,
                                        rmm::cuda_stream_view stream,
                                        rmm::mr::device_memory_resource* mr)
 {
@@ -122,11 +122,11 @@ std::unique_ptr<column> compute_column(table_view const table,
   // Execute the kernel
   auto table_device = table_device_view::create(table, stream);
   if (has_nulls) {
-    cudf::ast::detail::compute_column_kernel<MAX_BLOCK_SIZE, true>
+    cudf::detail::compute_column_kernel<MAX_BLOCK_SIZE, true>
       <<<config.num_blocks, config.num_threads_per_block, shmem_per_block, stream.value()>>>(
         *table_device, device_expression_data, *mutable_output_device);
   } else {
-    cudf::ast::detail::compute_column_kernel<MAX_BLOCK_SIZE, false>
+    cudf::detail::compute_column_kernel<MAX_BLOCK_SIZE, false>
       <<<config.num_blocks, config.num_threads_per_block, shmem_per_block, stream.value()>>>(
         *table_device, device_expression_data, *mutable_output_device);
   }
@@ -137,13 +137,11 @@ std::unique_ptr<column> compute_column(table_view const table,
 }  // namespace detail
 
 std::unique_ptr<column> compute_column(table_view const table,
-                                       expression const& expr,
+                                       ast::expression const& expr,
                                        rmm::mr::device_memory_resource* mr)
 {
   CUDF_FUNC_RANGE();
   return detail::compute_column(table, expr, rmm::cuda_stream_default, mr);
 }
 
-}  // namespace ast
-
 }  // namespace cudf