diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 193aa94409a7..ca1d96379a45 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -1034,7 +1034,6 @@ add_library( src/stream_compaction/distinct_helpers_nested_nan_unequal_ordered.cu src/stream_compaction/drop_nans.cu src/stream_compaction/drop_nulls.cu - src/stream_compaction/filter/filter.cu src/stream_compaction/stable_distinct.cu src/stream_compaction/unique.cu src/strings/attributes.cu diff --git a/cpp/benchmarks/filter/minmax_filter.cpp b/cpp/benchmarks/filter/minmax_filter.cpp index 512caa5fd9de..1c272d64c057 100644 --- a/cpp/benchmarks/filter/minmax_filter.cpp +++ b/cpp/benchmarks/filter/minmax_filter.cpp @@ -19,6 +19,7 @@ #include #include +#include #include #include @@ -120,31 +121,36 @@ void BM_filter_min_max(nvbench::state& state) auto const mem_stats_logger = cudf::memory_stats_logger(); state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { - auto stream = launch.get_stream().get_stream(); - auto mr = cudf::get_current_device_resource_ref(); + auto stream = launch.get_stream().get_stream(); + auto mr = cudf::get_current_device_resource_ref(); + auto filter_table = cudf::table_view{filter_column_views}; switch (engine) { case engine_type::AST: { auto predicate_table = cudf::table_view{{predicate_column->view()}}; - auto filter_table = cudf::table_view{filter_column_views}; auto const filter_boolean = cudf::compute_column(predicate_table, tree.back(), stream, mr); auto const result = cudf::apply_retention_mask(filter_table, filter_boolean->view(), stream, mr); } break; case engine_type::JIT: { - cudf::filter_input predicate_inputs[] = { + cudf::transform_input predicate_inputs[] = { predicate_column->view(), cudf::scalar_column_view(min_scalar_column->view()), cudf::scalar_column_view(max_scalar_column->view())}; - auto result = cudf::filter_extended(predicate_inputs, - udf, - filter_column_views, - cudf::udf_source_type::CUDA, - std::nullopt, - cudf::null_aware::NO, - cudf::output_nullability::PRESERVE, - stream, - mr); + cudf::transform_output outputs[] = {cudf::transform_output{ + cudf::data_type{cudf::type_to_id()}, cudf::output_nullability::PRESERVE}}; + auto const filter_boolean = cudf::transform(udf, + cudf::udf_source_type::CUDA, + cudf::null_aware::NO, + std::nullopt, + predicate_inputs, + outputs, + {}, + std::nullopt, + stream, + mr); + auto const result = + cudf::apply_retention_mask(filter_table, filter_boolean->view().column(0), stream, mr); } break; default: CUDF_UNREACHABLE("Unrecognised engine type requested"); } diff --git a/cpp/include/cudf/stream_compaction.hpp b/cpp/include/cudf/stream_compaction.hpp index 2b6a2e435914..e9a32b90f3ec 100644 --- a/cpp/include/cudf/stream_compaction.hpp +++ b/cpp/include/cudf/stream_compaction.hpp @@ -6,15 +6,12 @@ #pragma once #include -#include #include #include #include #include #include -#include -#include #include /** @@ -28,10 +25,6 @@ namespace CUDF_EXPORT cudf { * @{ */ -namespace ast { -struct expression; -} - /** * @brief Filters a table to remove null elements with threshold count. * @@ -409,128 +402,5 @@ std::unique_ptr stable_distinct( cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); -/** - * @brief Creates a new column by applying a filter function against every - * element of the input columns. - * - * Null values in the input columns are considered as not matching the filter. - * - * Computes: - * `out[i]... = predicate(columns[i]... ) ? (columns[i]...): not-applied`. - * - * Note that for every scalar in `columns` (columns of size 1), `columns[i] == - * input[0]` - * - * - * @throws std::invalid_argument if any of the input columns have different sizes (except scalars of - * size 1) - * @throws std::invalid_argument if the output or any of the inputs are not fixed-width or string - * types - * @throws cudf::logic_error if JIT is not supported by the runtime - * @throws std::invalid_argument if the size of `copy_mask` does not match the number of input - * columns - * - * The size of the resulting column is the size of the largest column. - * - * @param predicate_columns Immutable views of the predicate columns - * @param predicate_udf The PTX/CUDA string of the transform function to apply - * @param filter_columns Immutable view of the columns to be filtered - * @param is_ptx true: the UDF is treated as PTX code; false: the UDF is treated as CUDA code - * @param user_data User-defined device data to pass to the UDF. - * @param is_null_aware Signifies the UDF will receive row inputs as optional values - * @param predicate_nullability Specifies the nullability of the predicate output - * @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 - * @return The filtered target columns - */ -[[deprecated("Use filter_extended instead")]] std::vector> filter( - std::vector const& predicate_columns, - std::string const& predicate_udf, - std::vector const& filter_columns, - bool is_ptx, - std::optional user_data = std::nullopt, - null_aware is_null_aware = null_aware::NO, - output_nullability predicate_nullability = output_nullability::PRESERVE, - cuda::stream_ref stream = cudf::get_default_stream(), - rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); - -/** - * @brief Typedef for inputs to the filter function. Each input can be either a column or a - * scalar column. - */ -using filter_input = std::variant; - -/** - * @brief Creates a new column by applying a filter function against every - * element of the input columns. - * - * Null values in the input columns are considered as not matching the filter. - * - * Computes: - * `out[i]... = predicate(columns[i]... ) ? (columns[i]...): not-applied`. - * - * Note that for every scalar in `columns` (columns of size 1), `columns[i] == - * input[0]` - * - * - * @throws std::invalid_argument if any of the input columns have different sizes (except scalars of - * size 1) - * @throws std::invalid_argument if the output or any of the inputs are not fixed-width or string - * types - * @throws cudf::logic_error if JIT is not supported by the runtime - * @throws std::invalid_argument if the size of `copy_mask` does not match the number of input - * columns - * - * The size of the resulting column is the size of the largest column. - * - * @param predicate_inputs Immutable views of the predicate inputs (columns and scalars) - * @param predicate_udf The PTX/CUDA string of the transform function to apply - * @param filter_columns Immutable view of the columns to be filtered - * @param source_type The source type of the UDF - * @param user_data User-defined device data to pass to the UDF. - * @param is_null_aware Signifies the UDF will receive row inputs as optional values - * @param predicate_nullability Specifies the nullability of the predicate output - * @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 - * @return The filtered target columns - */ -std::vector> filter_extended( - std::span const> predicate_inputs, - std::string const& predicate_udf, - std::vector const& filter_columns, - cudf::udf_source_type source_type, - std::optional user_data = std::nullopt, - null_aware is_null_aware = null_aware::NO, - output_nullability predicate_nullability = output_nullability::PRESERVE, - cuda::stream_ref stream = cudf::get_default_stream(), - rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); - -/** - * @brief Creates new table by applying a filter function against every - * element of the input columns. - * - * Null values in the input columns are considered as not matching the filter. - * - * Computes: - * `out[i]... = predicate(columns[i]... ) ? (columns[i]...): not-applied`. - * - * @throws std::invalid_argument if the output or any of the inputs are not fixed-width or string - * types - * @throws cudf::logic_error if JIT is not supported by the runtime - * - * @param predicate_table The table used for predicate expression evaluation - * @param predicate_expr The predicate filter expression - * @param filter_table The table to be filtered - * @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 - * @return The filtered table - */ -std::unique_ptr
filter( - table_view const& predicate_table, - ast::expression const& predicate_expr, - table_view const& filter_table, - cuda::stream_ref stream = cudf::get_default_stream(), - rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); - /** @} */ } // namespace CUDF_EXPORT cudf diff --git a/cpp/src/io/parquet/reader_impl.cpp b/cpp/src/io/parquet/reader_impl.cpp index d263ac46d9a2..1a9180394d88 100644 --- a/cpp/src/io/parquet/reader_impl.cpp +++ b/cpp/src/io/parquet/reader_impl.cpp @@ -1023,22 +1023,22 @@ table_with_metadata reader_impl::finalize_output(read_mode mode, out_metadata.num_rows_per_source.clear(); bool use_jit = cudf::get_context().use_jit() || _options.use_jit_filter; + std::unique_ptr predicate; if (!use_jit) { - auto predicate = + predicate = cudf::detail::compute_column(*read_table, final_filter_expr.value().get(), _stream, _mr); - CUDF_EXPECTS(predicate->view().type().id() == type_id::BOOL8, - "Predicate filter should return a boolean"); - // Exclude columns present in filter only in output - auto output_table = cudf::detail::apply_mask( - only_output, *predicate, cudf::detail::mask_type::RETENTION, _stream, _mr); - return {encode_output_dict_columns(std::move(output_table)), std::move(out_metadata)}; } else { - auto output_table = cudf::filter( - read_table->view(), final_filter_expr.value().get(), only_output, _stream, _mr); - - return {encode_output_dict_columns(std::move(output_table)), std::move(out_metadata)}; + predicate = + cudf::compute_column_jit(*read_table, final_filter_expr.value().get(), _stream, _mr); } + + CUDF_EXPECTS(predicate->view().type().id() == type_id::BOOL8, + "Predicate filter should return a boolean"); + // Exclude columns present in filter only in output + auto output_table = cudf::detail::apply_mask( + only_output, *predicate, cudf::detail::mask_type::RETENTION, _stream, _mr); + return {encode_output_dict_columns(std::move(output_table)), std::move(out_metadata)}; } return {encode_output_dict_columns(std::make_unique
(std::move(out_columns))), std::move(out_metadata)}; diff --git a/cpp/src/stream_compaction/filter/filter.cu b/cpp/src/stream_compaction/filter/filter.cu deleted file mode 100644 index 67cac939a68f..000000000000 --- a/cpp/src/stream_compaction/filter/filter.cu +++ /dev/null @@ -1,159 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. - * SPDX-License-Identifier: Apache-2.0 - */ - -#include "jit/row_ir.hpp" - -#include -#include -#include -#include -#include -#include - -#include - -#include - -#include -#include - -namespace cudf { - -namespace detail { -std::unique_ptr
filter(std::string const& predicate_udf, - cudf::udf_source_type source_type, - null_aware is_null_aware, - std::optional user_data, - std::span predicate_inputs, - table_view const& filter_table, - output_nullability predicate_nullability, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) -{ - CUDF_EXPECTS(filter_table.num_columns() > 0, - "At least one column must be provided to filter.", - std::invalid_argument); - auto row_size = filter_table.num_rows(); - CUDF_EXPECTS(std::all_of(filter_table.begin(), - filter_table.end(), - [&](auto const& col) { return col.size() == row_size; }), - "All columns to filter must have the same number of rows.", - std::invalid_argument); - CUDF_EXPECTS(std::all_of(predicate_inputs.begin(), - predicate_inputs.end(), - [&](auto& input) { - if (auto* col = std::get_if(&input)) { - return col->size() == row_size; - } - return true; - }), - "All predicate input columns must have the same number of rows as the filter table.", - std::invalid_argument); - - transform_output outputs[] = {transform_output{data_type{type_id::BOOL8}, predicate_nullability}}; - - auto result = cudf::transform(predicate_udf, - source_type, - is_null_aware, - user_data, - predicate_inputs, - outputs, - {}, - filter_table.num_rows(), - stream, - mr); - - return apply_mask(filter_table, result->get_column(0), mask_type::RETENTION, stream, mr); -} - -} // namespace detail - -std::unique_ptr
filter(table_view const& predicate_table, - ast::expression const& predicate_expr, - table_view const& filter_table, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) -{ - auto args = cudf::detail::row_ir::ast_converter::filter(cudf::detail::row_ir::target::CUDA, - predicate_expr, - predicate_table, - {}, - "filter_operation", - stream, - mr); - - return detail::filter(args.udf, - args.source_type, - args.is_null_aware, - args.user_data, - args.inputs, - filter_table, - args.outputs[0].nullability, - stream, - mr); -} - -std::vector> filter_extended( - std::span const> predicate_inputs, - std::string const& predicate_udf, - std::vector const& filter_columns, - cudf::udf_source_type source_type, - std::optional user_data, - null_aware is_null_aware, - output_nullability predicate_nullability, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) -{ - CUDF_FUNC_RANGE(); - auto table = detail::filter(predicate_udf, - source_type, - is_null_aware, - user_data, - predicate_inputs, - table_view{filter_columns}, - predicate_nullability, - stream, - mr); - return table->release(); -} - -std::vector> filter(std::vector const& predicate_columns, - std::string const& predicate_udf, - std::vector const& filter_columns, - bool is_ptx, - std::optional user_data, - null_aware is_null_aware, - output_nullability predicate_nullability, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) -{ - // legacy behavior was to detect which column were scalars based on their sizes - std::vector> inputs; - -#pragma GCC diagnostic push -#pragma GCC diagnostic ignored "-Wdeprecated-declarations" - auto base_column = jit::get_transform_base_column(predicate_columns); - for (auto const& col : predicate_columns) { - if (jit::is_scalar(base_column->size(), col.size())) { -#pragma GCC diagnostic pop - inputs.emplace_back(scalar_column_view{col}); - } else { - inputs.emplace_back(col); - } - } - - auto table = detail::filter(predicate_udf, - is_ptx ? cudf::udf_source_type::PTX : cudf::udf_source_type::CUDA, - is_null_aware, - user_data, - inputs, - table_view{filter_columns}, - predicate_nullability, - stream, - mr); - return table->release(); -} - -} // namespace cudf diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 9153e1cbb82d..055f14533781 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -294,10 +294,6 @@ ConfigureTest( transform/one_hot_encode_tests.cpp ) -# ################################################################################################## -# * filter tests ------------------------------------------------------------------------- -ConfigureTest(FILTER_TEST filter/filter_test.cpp) - # ################################################################################################## # * interop tests ------------------------------------------------------------------------- ConfigureTest( diff --git a/cpp/tests/filter/filter_test.cpp b/cpp/tests/filter/filter_test.cpp deleted file mode 100644 index e32ca4f13c47..000000000000 --- a/cpp/tests/filter/filter_test.cpp +++ /dev/null @@ -1,406 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. - * SPDX-License-Identifier: Apache-2.0 - */ - -#include -#include -#include -#include -#include -#include - -#include -#include -#include - -namespace filters { - -struct FilterTestFixture : public cudf::test::BaseFixture { - protected: - static constexpr char const* udf = - R"***( - template - __device__ void is_equal(bool * out, T a, T b) { *out = (a == b); } - )***"; -}; - -template -struct FilterNumericTest : public FilterTestFixture {}; - -using NumericTypesNotBool = - cudf::test::Concat; - -TYPED_TEST_SUITE(FilterNumericTest, NumericTypesNotBool); - -TYPED_TEST(FilterNumericTest, NoAssertions) -{ - using T = TypeParam; - - auto a = cudf::test::fixed_width_column_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, - {1, 0, 1, 1, 1, 1, 1, 1, 0, 0}}; - auto b = cudf::test::fixed_width_column_wrapper{{0, 1, 2, 3, 8, 5, 6, 7, 4, 9}, - {0, 0, 1, 1, 1, 1, 1, 1, 0, 0}}; - - auto expected = cudf::test::fixed_width_column_wrapper{{2, 3, 5, 6, 7}, {1, 1, 1, 1, 1}}; - - std::vector> results; - - cudf::filter_input inputs[] = {a, b}; - - EXPECT_NO_THROW( - results = cudf::filter_extended( - inputs, this->udf, {a}, cudf::udf_source_type::CUDA, std::nullopt, cudf::null_aware::NO)); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, results[0]->view()); -} - -template -struct FilterChronoTest : public FilterTestFixture {}; - -TYPED_TEST_SUITE(FilterChronoTest, cudf::test::ChronoTypes); - -TYPED_TEST(FilterChronoTest, NoAssertions) -{ - using T = TypeParam; - - auto a = cudf::test::fixed_width_column_wrapper{ - {T{}, T{}, T{}, T{}, T{}, T{}, T{}, T{}, T{}, T{}}, {1, 0, 1, 1, 1, 1, 1, 1, 0, 0}}; - auto b = cudf::test::fixed_width_column_wrapper{ - {T{}, T{}, T{}, T{}, T{}, T{}, T{}, T{}, T{}, T{}}, {0, 0, 1, 1, 1, 1, 1, 1, 0, 0}}; - - auto expected = - cudf::test::fixed_width_column_wrapper{{T{}, T{}, T{}, T{}, T{}, T{}}, {1, 1, 1, 1, 1, 1}}; - - std::vector> results; - cudf::filter_input inputs[] = {a, b}; - EXPECT_NO_THROW( - results = cudf::filter_extended( - inputs, this->udf, {a}, cudf::udf_source_type::CUDA, std::nullopt, cudf::null_aware::NO)); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, results[0]->view()); -} - -template -struct FilterFixedPointTest : public FilterTestFixture {}; - -TYPED_TEST_SUITE(FilterFixedPointTest, cudf::test::FixedPointTypes); - -TYPED_TEST(FilterFixedPointTest, NoAssertions) -{ - using T = TypeParam; - - auto a = cudf::test::fixed_point_column_wrapper{ - {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, {1, 0, 1, 1, 1, 1, 1, 1, 0, 0}, numeric::scale_type{0}}; - auto b = cudf::test::fixed_point_column_wrapper{ - {0, 1, 2, 3, 8, 5, 6, 7, 4, 9}, {0, 0, 1, 1, 1, 1, 1, 1, 0, 0}, numeric::scale_type{0}}; - - auto expected = cudf::test::fixed_point_column_wrapper{ - {2, 3, 5, 6, 7}, {1, 1, 1, 1, 1}, numeric::scale_type{0}}; - - std::vector> results; - cudf::filter_input inputs[] = {a, b}; - - EXPECT_NO_THROW( - results = cudf::filter_extended( - inputs, this->udf, {a}, cudf::udf_source_type::CUDA, std::nullopt, cudf::null_aware::NO)); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, results[0]->view()); -} - -TEST_F(FilterTestFixture, StringNoAssertions) -{ - auto a = cudf::test::strings_column_wrapper{{"0", "1", "2", "3", "4", "5", "6", "7", "8", "9"}, - {1, 0, 1, 1, 1, 1, 1, 1, 0, 0}}; - auto b = cudf::test::strings_column_wrapper{{"0", "1", "2", "3", "8", "5", "6", "7", "8", "9"}, - {0, 0, 1, 1, 1, 1, 1, 1, 0, 0}}; - - auto expected = cudf::test::strings_column_wrapper{{"2", "3", "5", "6", "7"}, {1, 1, 1, 1, 1}}; - - std::vector> results; - cudf::filter_input inputs[] = {a, b}; - EXPECT_NO_THROW( - results = cudf::filter_extended( - inputs, this->udf, {a}, cudf::udf_source_type::CUDA, std::nullopt, cudf::null_aware::NO)); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, results[0]->view()); -} - -struct FilterAssertsTest : public FilterTestFixture {}; - -struct FilterTest : public FilterTestFixture {}; - -TEST_F(FilterTest, Basic) -{ - auto a = cudf::test::fixed_width_column_wrapper({1, 2, 3, 4, 5, 6, 7, 8, 9, 10}, - {1, 1, 1, 1, 1, 1, 1, 0, 0, 0}); - std::string cuda = R"***( -__device__ void is_even(bool* out, int32_t a) { *out = (a % 2 == 0); } - )***"; - - cudf::filter_input inputs[] = {a}; - - auto result = cudf::filter_extended( - inputs, cuda, {a}, cudf::udf_source_type::CUDA, std::nullopt, cudf::null_aware::NO); - auto expected = cudf::test::fixed_width_column_wrapper{{2, 4, 6}, {1, 1, 1}}; - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result[0]->view()); - - std::string null_cuda = R"***( -__device__ void is_even(cuda::std::optional* out, cuda::std::optional a) { *out = a.has_value() && (*a % 2 == 0); } - )***"; - - auto null_result = cudf::filter_extended( - inputs, null_cuda, {a}, cudf::udf_source_type::CUDA, std::nullopt, cudf::null_aware::YES); - auto null_expected = cudf::test::fixed_width_column_wrapper{{2, 4, 6}, {1, 1, 1}}; - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(null_expected, null_result[0]->view()); -} - -TEST_F(FilterTest, ScalarFilter) -{ - auto a = cudf::test::fixed_width_column_wrapper{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; - auto b = cudf::test::fixed_width_column_wrapper{2}; - std::string cuda = R"***( -__device__ void is_divisible(bool* out, int32_t a, int32_t b) { *out = ((a % b) == 0); } - )***"; - auto expected = cudf::test::fixed_width_column_wrapper{{2, 4, 6, 8, 10}}; - - cudf::filter_input inputs[] = {a, cudf::scalar_column_view(b)}; - - auto result = cudf::filter_extended( - inputs, cuda, {a}, cudf::udf_source_type::CUDA, std::nullopt, cudf::null_aware::NO); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result[0]->view()); -} - -TEST_F(FilterTest, MixedTypes) -{ - auto countries = cudf::test::strings_column_wrapper{ - "USA", "Canada", "Mexico", "Brazil", "Argentina", "France", "Germany", "Italy", "Spain"}; - auto average_tmp = cudf::test::fixed_width_column_wrapper{0, 0, 1, 1, 1, 1, 1, 1, 1}; - auto average_humidity = cudf::test::fixed_width_column_wrapper{0, 0, 0, 0, 1, 1, 1, 1, 1}; - auto timezones = cudf::test::strings_column_wrapper{ - "EST", "EST", "CST", "BRT", "ART", "CET", "CET", "CET", "CET"}; - - std::string cuda = R"***( -__device__ void filter(bool* out, - [[maybe_unused]] cudf::string_view country, - cudf::string_view tz, - float tmp, - float hum, - float min_tmp, - float max_tmp, - float min_hum, - float max_hum, - cudf::string_view tz1, - cudf::string_view tz2) -{ - *out = (tmp >= min_tmp && tmp <= max_tmp) && (hum >= min_hum && hum <= max_hum) && - (tz == tz1 || tz == tz2); -} - )***"; - - auto min_tmp = cudf::test::fixed_width_column_wrapper{0.5}; - auto max_tmp = cudf::test::fixed_width_column_wrapper{1}; - auto min_hum = cudf::test::fixed_width_column_wrapper{0.5}; - auto max_hum = cudf::test::fixed_width_column_wrapper{1}; - auto timezone1 = cudf::test::strings_column_wrapper{"CET"}; - auto timezone2 = cudf::test::strings_column_wrapper{"EST"}; - - cudf::filter_input inputs[] = {countries, - timezones, - average_tmp, - average_humidity, - cudf::scalar_column_view(min_tmp), - cudf::scalar_column_view(max_tmp), - cudf::scalar_column_view(min_hum), - cudf::scalar_column_view(max_hum), - cudf::scalar_column_view(timezone1), - cudf::scalar_column_view(timezone2)}; - - auto result = cudf::filter_extended(inputs, - cuda, - {countries, timezones}, - cudf::udf_source_type::CUDA, - std::nullopt, - cudf::null_aware::NO); - - EXPECT_EQ(result.size(), 2); - - auto expected_countries = - cudf::test::strings_column_wrapper{"France", "Germany", "Italy", "Spain"}; - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_countries, result[0]->view()); - - auto expected_timezones = cudf::test::strings_column_wrapper{"CET", "CET", "CET", "CET"}; - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_timezones, result[1]->view()); -} - -TEST_F(FilterTest, NullableMixedTypes) -{ - auto countries = cudf::test::strings_column_wrapper( - {"USA", "Canada", "Mexico", "Brazil", "Argentina", "France", "Germany", "Italy", "Spain"}, - {1, 1, 1, 1, 1, 0, 1, 0, 1}); - auto average_tmp = cudf::test::fixed_width_column_wrapper{0, 0, 1, 1, 1, 1, 1, 1, 1}; - auto average_humidity = cudf::test::fixed_width_column_wrapper{0, 0, 0, 0, 1, 1, 1, 1, 1}; - auto timezones = cudf::test::strings_column_wrapper{ - "EST", "EST", "CST", "BRT", "ART", "CET", "CET", "CET", "CET"}; - - std::string cuda = R"***( -__device__ void filter(bool* out, - [[maybe_unused]] cudf::string_view country, - cudf::string_view tz, - float tmp, - float hum, - float min_tmp, - float max_tmp, - float min_hum, - float max_hum, - cudf::string_view tz1, - cudf::string_view tz2) -{ - *out = (tmp >= min_tmp && tmp <= max_tmp) && (hum >= min_hum && hum <= max_hum) && - (tz == tz1 || tz == tz2); -} -)***"; - - auto min_tmp = cudf::test::fixed_width_column_wrapper{0.5}; - auto max_tmp = cudf::test::fixed_width_column_wrapper{1}; - auto min_hum = cudf::test::fixed_width_column_wrapper{0.5}; - auto max_hum = cudf::test::fixed_width_column_wrapper{1}; - auto timezone1 = cudf::test::strings_column_wrapper{"CET"}; - auto timezone2 = cudf::test::strings_column_wrapper{"EST"}; - - cudf::filter_input inputs[] = {countries, - timezones, - average_tmp, - average_humidity, - cudf::scalar_column_view(min_tmp), - cudf::scalar_column_view(max_tmp), - cudf::scalar_column_view(min_hum), - cudf::scalar_column_view(max_hum), - cudf::scalar_column_view(timezone1), - cudf::scalar_column_view(timezone2)}; - - auto result = cudf::filter_extended(inputs, - cuda, - {countries, timezones}, - cudf::udf_source_type::CUDA, - std::nullopt, - cudf::null_aware::NO); - - auto expected_countries = cudf::test::strings_column_wrapper({"Germany", "Spain"}, {true, true}); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_countries, result[0]->view()); - - auto expected_timezones = cudf::test::strings_column_wrapper({"CET", "CET"}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_timezones, result[1]->view()); -} - -struct ast_expression_executor { - static std::unique_ptr filter(cudf::ast::expression const& expr, - cudf::table_view const& table) - { - auto booleans = cudf::compute_column(table, expr); - return cudf::apply_retention_mask(table, booleans->view()); - } -}; - -struct jit_expression_executor { - static std::unique_ptr filter(cudf::ast::expression const& expr, - cudf::table_view const& table) - { - return cudf::filter(table, expr, table); - } -}; - -template -struct FilterExpressionTest : public cudf::test::BaseFixture { - std::unique_ptr a = - cudf::test::fixed_width_column_wrapper{{1, 2, 3, 4, 5, 6, 7, 8}, - {1, 1, 1, 1, 1, 1, 1, 0}} - .release(); - std::unique_ptr b = - cudf::test::fixed_width_column_wrapper{{1, 8, 3, 4, 5, 6, 7, 8}, - {1, 1, 1, 1, 1, 1, 1, 0}} - .release(); - std::unique_ptr bool_a = - cudf::test::fixed_width_column_wrapper{ - {false, false, true, true, false, false, true, true}, {1, 1, 1, 1, 1, 0, 0, 0}} - .release(); - std::unique_ptr bool_b = - cudf::test::fixed_width_column_wrapper{ - {false, true, false, true, true, true, false, true}, {1, 1, 1, 1, 0, 1, 1, 0}} - .release(); - - cudf::table_view table = cudf::table_view({a->view(), b->view(), bool_a->view(), bool_b->view()}); -}; - -using Executors = cudf::test::Types; - -TYPED_TEST_SUITE(FilterExpressionTest, Executors); - -TYPED_TEST(FilterExpressionTest, IsNull) -{ - using Executor = TypeParam; - - auto tree = cudf::ast::tree(); - auto& ref_0 = tree.push(cudf::ast::column_reference(0)); - auto& is_null_expr = tree.push(cudf::ast::operation(cudf::ast::ast_operator::IS_NULL, ref_0)); - auto& filter_expr = tree.push(cudf::ast::operation(cudf::ast::ast_operator::NOT, is_null_expr)); - auto result = Executor::filter(filter_expr, this->table); - auto expected_filter = - cudf::test::fixed_width_column_wrapper{{true, true, true, true, true, true, true, false}}; - auto expected_table = cudf::apply_retention_mask(this->table, expected_filter); - CUDF_TEST_EXPECT_TABLES_EQUAL(expected_table->view(), result->view()); -} - -TYPED_TEST(FilterExpressionTest, NullEqual) -{ - using Executor = TypeParam; - - auto tree = cudf::ast::tree(); - auto& ref_0 = tree.push(cudf::ast::column_reference(0)); - auto& ref_1 = tree.push(cudf::ast::column_reference(1)); - auto& null_equal_expr = - tree.push(cudf::ast::operation(cudf::ast::ast_operator::NULL_EQUAL, ref_0, ref_1)); - auto result = Executor::filter(null_equal_expr, this->table); - auto expected_filter = - cudf::test::fixed_width_column_wrapper{{true, false, true, true, true, true, true, true}}; - auto expected_table = cudf::apply_retention_mask(this->table, expected_filter); - CUDF_TEST_EXPECT_TABLES_EQUAL(expected_table->view(), result->view()); -} - -TYPED_TEST(FilterExpressionTest, NullLogicalAnd) -{ - using Executor = TypeParam; - - auto tree = cudf::ast::tree(); - auto& ref_2 = tree.push(cudf::ast::column_reference(2)); - auto& ref_3 = tree.push(cudf::ast::column_reference(3)); - auto& and_expr = - tree.push(cudf::ast::operation(cudf::ast::ast_operator::NULL_LOGICAL_AND, ref_2, ref_3)); - auto result = Executor::filter(and_expr, this->table); - auto expected_filter = cudf::test::fixed_width_column_wrapper{ - {false, false, false, true, false, false, false, false}}; - auto expected_table = cudf::apply_retention_mask(this->table, expected_filter); - CUDF_TEST_EXPECT_TABLES_EQUAL(expected_table->view(), result->view()); -} - -TYPED_TEST(FilterExpressionTest, NullLogicalOr) -{ - using Executor = TypeParam; - - auto tree = cudf::ast::tree(); - auto& ref_2 = tree.push(cudf::ast::column_reference(2)); - auto& ref_3 = tree.push(cudf::ast::column_reference(3)); - auto& or_expr = - tree.push(cudf::ast::operation(cudf::ast::ast_operator::NULL_LOGICAL_OR, ref_2, ref_3)); - auto result = Executor::filter(or_expr, this->table); - auto expected_filter = cudf::test::fixed_width_column_wrapper{ - {false, true, true, true, false, true, false, false}}; - auto expected_table = cudf::apply_retention_mask(this->table, expected_filter); - CUDF_TEST_EXPECT_TABLES_EQUAL(expected_table->view(), result->view()); -} - -} // namespace filters - -CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/streams/stream_compaction_test.cpp b/cpp/tests/streams/stream_compaction_test.cpp index 5078abc18c7a..f86b420d2c42 100644 --- a/cpp/tests/streams/stream_compaction_test.cpp +++ b/cpp/tests/streams/stream_compaction_test.cpp @@ -9,7 +9,6 @@ #include #include -#include #include #include #include @@ -395,41 +394,6 @@ TEST_F(StreamCompactionTest, ApplyDeletionMask) CUDF_TEST_EXPECT_TABLES_EQUAL(expected, *result); } -TEST_F(StreamCompactionTest, FilterUDF) -{ - auto const col = int32s_col{0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14}; - auto col_ref_0 = cudf::ast::column_reference(0); - auto const expected = int32s_col{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}.release(); - cudf::filter_input inputs[] = {col}; - auto const result = cudf::filter_extended(inputs, - R"***( -__device__ void filter(bool * out, int32_t a){ - *out = a < 10; -})***", - {col}, - cudf::udf_source_type::CUDA, - std::nullopt, - cudf::null_aware::NO, - cudf::output_nullability::PRESERVE, - cudf::test::get_default_stream()); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*expected, *result[0]); -} - -TEST_F(StreamCompactionTest, FilterASTJit) -{ - auto const col = int32s_col{0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14}; - auto col_ref_0 = cudf::ast::column_reference(0); - auto max_scalar = cudf::numeric_scalar( - 10, true, cudf::test::get_default_stream(), cudf::get_current_device_resource_ref()); - auto const max_literal = cudf::ast::literal(max_scalar); - auto expression = cudf::ast::operation(cudf::ast::ast_operator::LESS, col_ref_0, max_literal); - cudf::table_view input({col}); - auto const col_expected = int32s_col{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}; - cudf::table_view expected({col_expected}); - auto const result = cudf::filter(input, expression, input, cudf::test::get_default_stream()); - CUDF_TEST_EXPECT_TABLES_EQUAL(expected, *result); -} - TEST_F(StreamCompactionTest, UniqueCountColumn) { std::vector const input = {1, 3, 3, 4, 31, 1, 8, 2, 0, 4, 1, diff --git a/python/pylibcudf/pylibcudf/libcudf/stream_compaction.pxd b/python/pylibcudf/pylibcudf/libcudf/stream_compaction.pxd index b3b85822c908..d59d80f3456b 100644 --- a/python/pylibcudf/pylibcudf/libcudf/stream_compaction.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/stream_compaction.pxd @@ -6,7 +6,6 @@ from libcpp.vector cimport vector from pylibcudf.exception_handler cimport libcudf_exception_handler from pylibcudf.libcudf.column.column cimport column from pylibcudf.libcudf.column.column_view cimport column_view -from pylibcudf.libcudf.expressions cimport expression from pylibcudf.libcudf.table.table cimport table from pylibcudf.libcudf.table.table_view cimport table_view from pylibcudf.libcudf.types cimport ( @@ -99,11 +98,3 @@ cdef extern from "cudf/stream_compaction.hpp" namespace "cudf" nogil: cudaStream_t stream, device_async_resource_ref mr ) except +libcudf_exception_handler - - cdef unique_ptr[table] filter( - table_view predicate_table, - const expression& predicate_expr, - table_view filter_table, - cudaStream_t stream, - device_async_resource_ref mr - ) except +libcudf_exception_handler diff --git a/python/pylibcudf/pylibcudf/stream_compaction.pxd b/python/pylibcudf/pylibcudf/stream_compaction.pxd index 705eea5e0a1f..5b37f3bce3e6 100644 --- a/python/pylibcudf/pylibcudf/stream_compaction.pxd +++ b/python/pylibcudf/pylibcudf/stream_compaction.pxd @@ -10,7 +10,6 @@ from pylibcudf.libcudf.types cimport ( from rmm.pylibrmm.memory_resource cimport DeviceMemoryResource from .column cimport Column -from .expressions cimport Expression from .table cimport Table @@ -88,11 +87,3 @@ cpdef Table stable_distinct( object stream = *, DeviceMemoryResource mr = *, ) - -cpdef Table filter( - Table predicate_table, - Expression predicate_expr, - Table filter_table, - object stream = *, - DeviceMemoryResource mr = *, -) diff --git a/python/pylibcudf/pylibcudf/stream_compaction.pyi b/python/pylibcudf/pylibcudf/stream_compaction.pyi index 617cfce17e4d..03a0ecc948b8 100644 --- a/python/pylibcudf/pylibcudf/stream_compaction.pyi +++ b/python/pylibcudf/pylibcudf/stream_compaction.pyi @@ -6,7 +6,6 @@ from enum import IntEnum from rmm.pylibrmm.memory_resource import DeviceMemoryResource from pylibcudf.column import Column -from pylibcudf.expressions import Expression from pylibcudf.table import Table from pylibcudf.types import NanEquality, NullEquality from pylibcudf.utils import CudaStreamLike @@ -83,10 +82,3 @@ def stable_distinct( stream: CudaStreamLike | None = None, mr: DeviceMemoryResource | None = None, ) -> Table: ... -def filter( - predicate_table: Table, - predicate_expr: Expression, - filter_table: Table, - stream: CudaStreamLike | None = None, - mr: DeviceMemoryResource | None = None, -) -> Table: ... diff --git a/python/pylibcudf/pylibcudf/stream_compaction.pyx b/python/pylibcudf/pylibcudf/stream_compaction.pyx index f7548c5ebcf0..5207bd39e0c2 100644 --- a/python/pylibcudf/pylibcudf/stream_compaction.pyx +++ b/python/pylibcudf/pylibcudf/stream_compaction.pyx @@ -1,7 +1,6 @@ # SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 -from cython.operator cimport dereference from libcpp.memory cimport unique_ptr from libcpp.utility cimport move from libcpp.vector cimport vector @@ -23,7 +22,6 @@ from rmm.pylibrmm.memory_resource cimport DeviceMemoryResource from rmm.pylibrmm.stream cimport Stream from .column cimport Column -from .expressions cimport Expression from .table cimport Table from .utils cimport _get_stream, _get_memory_resource @@ -43,7 +41,6 @@ __all__ = [ "distinct_indices", "drop_nans", "drop_nulls", - "filter", "stable_distinct", "unique", ] @@ -404,47 +401,4 @@ cpdef Table stable_distinct( return Table.from_libcudf(move(c_result), _stream, mr) -cpdef Table filter( - Table predicate_table, - Expression predicate_expr, - Table filter_table, - object stream: CudaStreamLike | None = None, - DeviceMemoryResource mr=None, -): - """Filters a table using a predicate expression. - - For details, see :cpp:func:`filter`. - - Parameters - ---------- - predicate_table : Table - The table used for predicate expression evaluation. - predicate_expr : Expression - The predicate filter expression. - filter_table : Table - The table to be filtered. - - Returns - ------- - Table - The filtered table. - """ - cdef unique_ptr[table] c_result - - cdef Stream _stream = _get_stream(stream) - cdef cudaStream_t _cs = _stream.view().value() - mr = _get_memory_resource(mr) - - cdef table_view c_predicate_table = predicate_table.view() - cdef table_view c_filter_table = filter_table.view() - with nogil: - c_result = cpp_stream_compaction.filter( - c_predicate_table, - dereference(predicate_expr.c_obj.get()), - c_filter_table, - _cs, - mr.get_mr() - ) - return Table.from_libcudf(move(c_result), _stream, mr) - DuplicateKeepOption.__str__ = DuplicateKeepOption.__repr__