Skip to content

Use cuda::std:: traits and utilities for AST operators #19179

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 7 commits into from
Jun 24, 2025
Merged

Use cuda::std:: traits and utilities for AST operators #19179

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
d5bce56
Use CCCL traits and utilities for AST operators
PointKernel Jun 16, 2025
0656507
Merge remote-tracking branch 'upstream/branch-25.08' into cccl-ast-op…
PointKernel Jun 17, 2025
6f75170
Move AST host APIs to a cpp header
PointKernel Jun 17, 2025
13198bd
Update copyright
PointKernel Jun 17, 2025
c147017
Cleanups
PointKernel Jun 18, 2025
bebe336
Merge remote-tracking branch 'upstream/branch-25.08' into cccl-ast-op…
PointKernel Jun 18, 2025
a5a2098
Merge branch 'branch-25.08' into cccl-ast-operators
PointKernel Jun 18, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions cpp/benchmarks/common/ndsh_data_generator/ndsh_data_generator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2024, NVIDIA CORPORATION.
* Copyright (c) 2024-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand All @@ -21,7 +21,7 @@

#include <cudf_test/column_wrapper.hpp>

#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/detail/operators.cuh>
#include <cudf/ast/expressions.hpp>
#include <cudf/binaryop.hpp>
#include <cudf/detail/nvtx/ranges.hpp>
Expand Down
2 changes: 1 addition & 1 deletion cpp/benchmarks/common/ndsh_data_generator/table_helpers.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@
#include "random_column_generator.hpp"

#include <cudf/aggregation.hpp>
#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/detail/operators.cuh>
#include <cudf/ast/expressions.hpp>
#include <cudf/binaryop.hpp>
#include <cudf/column/column_factories.hpp>
Expand Down
2 changes: 1 addition & 1 deletion cpp/include/cudf/ast/detail/expression_evaluator.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,7 @@
#pragma once

#include <cudf/ast/detail/expression_parser.hpp>
#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/detail/operators.cuh>
#include <cudf/ast/expressions.hpp>
#include <cudf/column/column_device_view.cuh>
#include <cudf/column/column_factories.hpp>
Expand Down
2 changes: 1 addition & 1 deletion cpp/include/cudf/ast/detail/expression_parser.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,7 @@
*/
#pragma once

#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/detail/operators.cuh>
#include <cudf/ast/expressions.hpp>
#include <cudf/detail/utilities/integer_utils.hpp>
#include <cudf/table/table_view.hpp>
Expand Down
240 changes: 119 additions & 121 deletions cpp/include/cudf/ast/detail/operators.hpp → cpp/include/cudf/ast/detail/operators.cuh
View file
Open in desktop
Copy link
Contributor

@bdice bdice Jun 16, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Does this need to be a .cuh? See rules in developer guide. This is primarily __device__ functions, so I can see the argument, but I believe CUDF_HOST_DEVICE is meant to be usable by host compilers too.

Copy link
Member Author

@PointKernel PointKernel Jun 16, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm not sure CUDF_HOST_DEVICE was ever meant to support host-only compilers. I had assumed it was just a temporary workaround from when CCCL wasn’t mature, and we had to rely on relaxed constexpr to use STL utilities on device. Back then, we didn’t really try to separate host and device usage because we didn’t have the option. Also, since CUDA rdc is disabled in libcudf, device function definitions can't be separated from declarations. If they contain host-device code, even headers ending in .hpp have likely always been compiled with nvcc.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Many of the pieces like ast_operator_dispatcher, ast_operator_return_type, ast_operator_arity, etc. are meant to be usable in host-only code, that's how we build up the AST program.

Copy link
Member Author

@PointKernel PointKernel Jun 17, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Let me move ast_operator_return_type and ast_operator_arity to operators.hpp. As for ast_operator_dispatcher, it seems that it's used in both host and device code so I will will probably leave it in the .cuh?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Does this need to be a .cuh? See rules in developer guide. This is primarily device functions, so I can see the argument, but I believe CUDF_HOST_DEVICE is meant to be usable by host compilers too.

Counterpoint: Given that the operator()()s run in a __device__ context, it appears that there isn't any getting around nvcc having to process this file. Wouldn't switching it to a .hpp then at best match the current compile time?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

To be clear, CUDF_HOST_DEVICE does not rely on relaxed constexpr. The whole point of specifying CUDF_HOST_DEVICE is so that we get valid host and device callables for those functions without any UB around constexpr functions.

Large diffs are not rendered by default.

2 changes: 1 addition & 1 deletion cpp/src/ast/expression_parser.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@
* limitations under the License.
*/
#include <cudf/ast/detail/expression_parser.hpp>
#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/detail/operators.cuh>
#include <cudf/ast/expressions.hpp>
#include <cudf/types.hpp>
#include <cudf/utilities/error.hpp>
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/ast/expressions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2021-2024, NVIDIA CORPORATION.
* Copyright (c) 2021-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand All @@ -15,7 +15,7 @@
*/
#include <cudf/ast/detail/expression_parser.hpp>
#include <cudf/ast/detail/expression_transformer.hpp>
#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/detail/operators.cuh>
#include <cudf/ast/expressions.hpp>
#include <cudf/types.hpp>
#include <cudf/utilities/error.hpp>
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/ast/operators.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/detail/operators.cuh>
#include <cudf/utilities/error.hpp>
#include <cudf/utilities/type_dispatcher.hpp>

Expand Down
2 changes: 1 addition & 1 deletion cpp/src/io/parquet/bloom_filter_reader.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@
#include "reader_impl_helpers.hpp"

#include <cudf/ast/detail/expression_transformer.hpp>
#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/detail/operators.cuh>
#include <cudf/ast/expressions.hpp>
#include <cudf/detail/cuco_helpers.hpp>
#include <cudf/detail/transform.hpp>
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/io/parquet/experimental/dictionary_page_filter.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@
#include "hybrid_scan_helpers.hpp"

#include <cudf/ast/detail/expression_transformer.hpp>
#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/detail/operators.cuh>
#include <cudf/ast/expressions.hpp>
#include <cudf/detail/cuco_helpers.hpp>
#include <cudf/detail/transform.hpp>
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/io/parquet/predicate_pushdown.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@
#include "stats_filter_helpers.hpp"

#include <cudf/ast/detail/expression_transformer.hpp>
#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/detail/operators.cuh>
#include <cudf/ast/expressions.hpp>
#include <cudf/column/column_factories.hpp>
#include <cudf/detail/iterator.cuh>
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/io/parquet/stats_filter_helpers.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@

#include "io/parquet/parquet_common.hpp"

#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/detail/operators.cuh>
#include <cudf/ast/expressions.hpp>
#include <cudf/utilities/error.hpp>
#include <cudf/utilities/traits.hpp>
Expand Down