From 9fae839bfa3b58ad76a6d3825fb9713c972a1b51 Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Tue, 18 May 2021 17:04:57 -0500 Subject: [PATCH 1/7] Do not add nulls to the hash table when null_equality::NOT_EQUAL is passed to left_anti_join and left_semi_join. Major performance gains. Added benchmarks. --- cpp/benchmarks/join/join_benchmark.cu | 137 ++++++++++++++++++++++++-- cpp/include/cudf/join.hpp | 4 +- cpp/src/join/semi_join.cu | 22 +++-- 3 files changed, 148 insertions(+), 15 deletions(-) diff --git a/cpp/benchmarks/join/join_benchmark.cu b/cpp/benchmarks/join/join_benchmark.cu index d1c11696ddd..824ae35fc65 100644 --- a/cpp/benchmarks/join/join_benchmark.cu +++ b/cpp/benchmarks/join/join_benchmark.cu @@ -37,8 +37,8 @@ template class Join : public cudf::benchmark { }; -template -static void BM_join(benchmark::State &state) +template +static void BM_join(Join JoinFunc, benchmark::State& state) { const cudf::size_type build_table_size{(cudf::size_type)state.range(0)}; const cudf::size_type probe_table_size{(cudf::size_type)state.range(1)}; @@ -105,20 +105,69 @@ static void BM_join(benchmark::State &state) for (auto _ : state) { cuda_event_timer raii(state, true, rmm::cuda_stream_default); - auto result = cudf::inner_join( + auto result = JoinFunc( probe_table, build_table, columns_to_join, columns_to_join, cudf::null_equality::UNEQUAL); } } -#define JOIN_BENCHMARK_DEFINE(name, key_type, payload_type, nullable) \ - BENCHMARK_TEMPLATE_DEFINE_F(Join, name, key_type, payload_type) \ - (::benchmark::State & st) { BM_join(st); } +#define JOIN_BENCHMARK_DEFINE(name, key_type, payload_type, nullable) \ + BENCHMARK_TEMPLATE_DEFINE_F(Join, name, key_type, payload_type) \ + (::benchmark::State & st) \ + { \ + auto join = [](cudf::table_view const& left, \ + cudf::table_view const& right, \ + std::vector const& left_on, \ + std::vector const& right_on, \ + cudf::null_equality compare_nulls) { \ + return cudf::inner_join(left, right, left_on, right_on, compare_nulls); \ + }; \ + BM_join(join, st); \ + } JOIN_BENCHMARK_DEFINE(join_32bit, int32_t, int32_t, false); JOIN_BENCHMARK_DEFINE(join_64bit, int64_t, int64_t, false); JOIN_BENCHMARK_DEFINE(join_32bit_nulls, int32_t, int32_t, true); JOIN_BENCHMARK_DEFINE(join_64bit_nulls, int64_t, int64_t, true); +#define LEFT_ANTI_JOIN_BENCHMARK_DEFINE(name, key_type, payload_type, nullable) \ + BENCHMARK_TEMPLATE_DEFINE_F(Join, name, key_type, payload_type) \ + (::benchmark::State & st) \ + { \ + auto join = [](cudf::table_view const& left, \ + cudf::table_view const& right, \ + std::vector const& left_on, \ + std::vector const& right_on, \ + cudf::null_equality compare_nulls) { \ + return cudf::left_anti_join(left, right, left_on, right_on, compare_nulls); \ + }; \ + BM_join(join, st); \ + } + +LEFT_ANTI_JOIN_BENCHMARK_DEFINE(left_anti_join_32bit, int32_t, int32_t, false); +LEFT_ANTI_JOIN_BENCHMARK_DEFINE(left_anti_join_64bit, int64_t, int64_t, false); +LEFT_ANTI_JOIN_BENCHMARK_DEFINE(left_anti_join_32bit_nulls, int32_t, int32_t, true); +LEFT_ANTI_JOIN_BENCHMARK_DEFINE(left_anti_join_64bit_nulls, int64_t, int64_t, true); + +#define LEFT_SEMI_JOIN_BENCHMARK_DEFINE(name, key_type, payload_type, nullable) \ + BENCHMARK_TEMPLATE_DEFINE_F(Join, name, key_type, payload_type) \ + (::benchmark::State & st) \ + { \ + auto join = [](cudf::table_view const& left, \ + cudf::table_view const& right, \ + std::vector const& left_on, \ + std::vector const& right_on, \ + cudf::null_equality compare_nulls) { \ + return cudf::left_semi_join(left, right, left_on, right_on, compare_nulls); \ + }; \ + BM_join(join, st); \ + } + +LEFT_SEMI_JOIN_BENCHMARK_DEFINE(left_semi_join_32bit, int32_t, int32_t, false); +LEFT_SEMI_JOIN_BENCHMARK_DEFINE(left_semi_join_64bit, int64_t, int64_t, false); +LEFT_SEMI_JOIN_BENCHMARK_DEFINE(left_semi_join_32bit_nulls, int32_t, int32_t, true); +LEFT_SEMI_JOIN_BENCHMARK_DEFINE(left_semi_join_64bit_nulls, int64_t, int64_t, true); + +// join ----------------------------------------------------------------------- BENCHMARK_REGISTER_F(Join, join_32bit) ->Unit(benchmark::kMillisecond) ->Args({100'000, 100'000}) @@ -144,7 +193,7 @@ BENCHMARK_REGISTER_F(Join, join_32bit_nulls) ->Args({100'000, 1'000'000}) ->Args({10'000'000, 10'000'000}) ->Args({10'000'000, 40'000'000}) - ->Args({10'000'000, 100'000'000}) + ->Args({10'000'00, 100'000'00}) ->Args({100'000'000, 100'000'000}) ->Args({80'000'000, 240'000'000}) ->UseManualTime(); @@ -154,3 +203,77 @@ BENCHMARK_REGISTER_F(Join, join_64bit_nulls) ->Args({50'000'000, 50'000'000}) ->Args({40'000'000, 120'000'000}) ->UseManualTime(); + +// left anti-join ------------------------------------------------------------- +BENCHMARK_REGISTER_F(Join, left_anti_join_32bit) + ->Unit(benchmark::kMillisecond) + ->Args({100'000, 100'000}) + ->Args({100'000, 400'000}) + ->Args({100'000, 1'000'000}) + ->Args({10'000'000, 10'000'000}) + ->Args({10'000'000, 40'000'000}) + ->Args({10'000'000, 100'000'000}) + ->Args({100'000'000, 100'000'000}) + ->Args({80'000'000, 240'000'000}) + ->UseManualTime(); + +BENCHMARK_REGISTER_F(Join, left_anti_join_64bit) + ->Unit(benchmark::kMillisecond) + ->Args({50'000'000, 50'000'000}) + ->Args({40'000'000, 120'000'000}) + ->UseManualTime(); + +BENCHMARK_REGISTER_F(Join, left_anti_join_32bit_nulls) + ->Unit(benchmark::kMillisecond) + ->Args({100'000, 100'000}) + ->Args({100'000, 400'000}) + ->Args({100'000, 1'000'000}) + ->Args({10'000'000, 10'000'000}) + ->Args({10'000'000, 40'000'000}) + ->Args({10'000'00, 100'000'00}) + ->Args({100'000'000, 100'000'000}) + ->Args({80'000'000, 240'000'000}) + ->UseManualTime(); + +BENCHMARK_REGISTER_F(Join, left_anti_join_64bit_nulls) + ->Unit(benchmark::kMillisecond) + ->Args({50'000'000, 50'000'000}) + ->Args({40'000'000, 120'000'000}) + ->UseManualTime(); + +// left semi-join ------------------------------------------------------------- +BENCHMARK_REGISTER_F(Join, left_semi_join_32bit) + ->Unit(benchmark::kMillisecond) + ->Args({100'000, 100'000}) + ->Args({100'000, 400'000}) + ->Args({100'000, 1'000'000}) + ->Args({10'000'000, 10'000'000}) + ->Args({10'000'000, 40'000'000}) + ->Args({10'000'000, 100'000'000}) + ->Args({100'000'000, 100'000'000}) + ->Args({80'000'000, 240'000'000}) + ->UseManualTime(); + +BENCHMARK_REGISTER_F(Join, left_semi_join_64bit) + ->Unit(benchmark::kMillisecond) + ->Args({50'000'000, 50'000'000}) + ->Args({40'000'000, 120'000'000}) + ->UseManualTime(); + +BENCHMARK_REGISTER_F(Join, left_semi_join_32bit_nulls) + ->Unit(benchmark::kMillisecond) + ->Args({100'000, 100'000}) + ->Args({100'000, 400'000}) + ->Args({100'000, 1'000'000}) + ->Args({10'000'000, 10'000'000}) + ->Args({10'000'000, 40'000'000}) + ->Args({10'000'00, 100'000'00}) + ->Args({100'000'000, 100'000'000}) + ->Args({80'000'000, 240'000'000}) + ->UseManualTime(); + +BENCHMARK_REGISTER_F(Join, left_semi_join_64bit_nulls) + ->Unit(benchmark::kMillisecond) + ->Args({50'000'000, 50'000'000}) + ->Args({40'000'000, 120'000'000}) + ->UseManualTime(); \ No newline at end of file diff --git a/cpp/include/cudf/join.hpp b/cpp/include/cudf/join.hpp index 5a2c913d4c3..428a4195bf8 100644 --- a/cpp/include/cudf/join.hpp +++ b/cpp/include/cudf/join.hpp @@ -424,13 +424,13 @@ std::unique_ptr> left_anti_join( * TableB: {{1, 2, 3}, {1, 2, 5}} * left_on: {0} * right_on: {1} - * Result: {{0}, {1}} + * Result: {{0}} * * TableA: {{0, 1, 2}, {1, 2, 5}} * TableB: {{1, 2, 3}} * left_on: {0} * right_on: {0} - * Result: { {0} {1} } + * Result: { {0}, {1} } * @endcode * * @throw cudf::logic_error if number of elements in `left_on` or `right_on` diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index c029f0272da..5fc3a3dbcda 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -92,12 +92,22 @@ std::unique_ptr> left_semi_anti_join( equality_build); auto hash_table = *hash_table_ptr; - thrust::for_each_n(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - right_num_rows, - [hash_table] __device__(size_type idx) mutable { - hash_table.insert(thrust::make_pair(idx, true)); - }); + // if compare_nulls == NOT_EQUAL, we can simply ignore any rows that are + // entirely NULL as they will never compare to equal. + auto const row_bitmask = (compare_nulls == null_equality::EQUAL) + ? rmm::device_buffer{0, stream} + : cudf::detail::bitmask_and(right_flattened_keys, stream); + // skip rows that are null here. + thrust::for_each_n( + rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + right_num_rows, + [hash_table, row_bitmask = static_cast(row_bitmask.data())] __device__( + size_type idx) mutable { + if (!row_bitmask || cudf::bit_is_set(row_bitmask, idx)) { + hash_table.insert(thrust::make_pair(idx, true)); + } + }); // // Now we have a hash table, we need to iterate over the rows of the left table From a4e3f309300546f1809ae8d497546aef56567783 Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Tue, 18 May 2021 17:32:23 -0500 Subject: [PATCH 2/7] Restore a couple of benchmark parameters that had been artificially lowered. --- cpp/benchmarks/join/join_benchmark.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/benchmarks/join/join_benchmark.cu b/cpp/benchmarks/join/join_benchmark.cu index 824ae35fc65..dce652add2b 100644 --- a/cpp/benchmarks/join/join_benchmark.cu +++ b/cpp/benchmarks/join/join_benchmark.cu @@ -193,7 +193,7 @@ BENCHMARK_REGISTER_F(Join, join_32bit_nulls) ->Args({100'000, 1'000'000}) ->Args({10'000'000, 10'000'000}) ->Args({10'000'000, 40'000'000}) - ->Args({10'000'00, 100'000'00}) + ->Args({10'000'000, 100'000'00}) ->Args({100'000'000, 100'000'000}) ->Args({80'000'000, 240'000'000}) ->UseManualTime(); @@ -230,7 +230,7 @@ BENCHMARK_REGISTER_F(Join, left_anti_join_32bit_nulls) ->Args({100'000, 1'000'000}) ->Args({10'000'000, 10'000'000}) ->Args({10'000'000, 40'000'000}) - ->Args({10'000'00, 100'000'00}) + ->Args({10'000'000, 100'000'00}) ->Args({100'000'000, 100'000'000}) ->Args({80'000'000, 240'000'000}) ->UseManualTime(); @@ -267,7 +267,7 @@ BENCHMARK_REGISTER_F(Join, left_semi_join_32bit_nulls) ->Args({100'000, 1'000'000}) ->Args({10'000'000, 10'000'000}) ->Args({10'000'000, 40'000'000}) - ->Args({10'000'00, 100'000'00}) + ->Args({10'000'000, 100'000'00}) ->Args({100'000'000, 100'000'000}) ->Args({80'000'000, 240'000'000}) ->UseManualTime(); From 97e1a4745fe5e3a74a1a46155235bfcace5a3587 Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Tue, 18 May 2021 17:33:37 -0500 Subject: [PATCH 3/7] More benchmark number tweaks. --- cpp/benchmarks/join/join_benchmark.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/benchmarks/join/join_benchmark.cu b/cpp/benchmarks/join/join_benchmark.cu index dce652add2b..0eb3989e480 100644 --- a/cpp/benchmarks/join/join_benchmark.cu +++ b/cpp/benchmarks/join/join_benchmark.cu @@ -193,7 +193,7 @@ BENCHMARK_REGISTER_F(Join, join_32bit_nulls) ->Args({100'000, 1'000'000}) ->Args({10'000'000, 10'000'000}) ->Args({10'000'000, 40'000'000}) - ->Args({10'000'000, 100'000'00}) + ->Args({10'000'000, 100'000'000}) ->Args({100'000'000, 100'000'000}) ->Args({80'000'000, 240'000'000}) ->UseManualTime(); @@ -230,7 +230,7 @@ BENCHMARK_REGISTER_F(Join, left_anti_join_32bit_nulls) ->Args({100'000, 1'000'000}) ->Args({10'000'000, 10'000'000}) ->Args({10'000'000, 40'000'000}) - ->Args({10'000'000, 100'000'00}) + ->Args({10'000'000, 100'000'000}) ->Args({100'000'000, 100'000'000}) ->Args({80'000'000, 240'000'000}) ->UseManualTime(); @@ -267,7 +267,7 @@ BENCHMARK_REGISTER_F(Join, left_semi_join_32bit_nulls) ->Args({100'000, 1'000'000}) ->Args({10'000'000, 10'000'000}) ->Args({10'000'000, 40'000'000}) - ->Args({10'000'000, 100'000'00}) + ->Args({10'000'000, 100'000'000}) ->Args({100'000'000, 100'000'000}) ->Args({80'000'000, 240'000'000}) ->UseManualTime(); From 99827d1e373185cc3685fec64cdef86e4ce1e3db Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Tue, 18 May 2021 17:35:36 -0500 Subject: [PATCH 4/7] Newline in benchmarks --- cpp/benchmarks/join/join_benchmark.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/benchmarks/join/join_benchmark.cu b/cpp/benchmarks/join/join_benchmark.cu index 0eb3989e480..357852629f4 100644 --- a/cpp/benchmarks/join/join_benchmark.cu +++ b/cpp/benchmarks/join/join_benchmark.cu @@ -276,4 +276,5 @@ BENCHMARK_REGISTER_F(Join, left_semi_join_64bit_nulls) ->Unit(benchmark::kMillisecond) ->Args({50'000'000, 50'000'000}) ->Args({40'000'000, 120'000'000}) - ->UseManualTime(); \ No newline at end of file + ->UseManualTime(); + \ No newline at end of file From 8226b0c4dd5ef346a1d101b75ec4b153605aab21 Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Thu, 20 May 2021 10:56:52 -0500 Subject: [PATCH 5/7] Add missing join tests: Semi-join with nulls UNEQUAL, and anti-join with both nulls EQUAL and UNEQUAL. --- cpp/tests/CMakeLists.txt | 2 +- cpp/tests/join/semi_anti_join_tests.cpp | 217 ++++++++++++++++++++++++ cpp/tests/join/semi_join_tests.cpp | 109 ------------ 3 files changed, 218 insertions(+), 110 deletions(-) create mode 100644 cpp/tests/join/semi_anti_join_tests.cpp delete mode 100644 cpp/tests/join/semi_join_tests.cpp diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 914edf95fd1..2726a6ce255 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -84,7 +84,7 @@ ConfigureTest(GROUPBY_TEST ConfigureTest(JOIN_TEST join/join_tests.cpp join/cross_join_tests.cpp - join/semi_join_tests.cpp) + join/semi_anti_join_tests.cpp) ################################################################################################### # - is_sorted tests ------------------------------------------------------------------------------- diff --git a/cpp/tests/join/semi_anti_join_tests.cpp b/cpp/tests/join/semi_anti_join_tests.cpp new file mode 100644 index 00000000000..5b38bafb122 --- /dev/null +++ b/cpp/tests/join/semi_anti_join_tests.cpp @@ -0,0 +1,217 @@ +/* + * Copyright (c) 2021, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +#include + +template +using column_wrapper = cudf::test::fixed_width_column_wrapper; +using strcol_wrapper = cudf::test::strings_column_wrapper; +using column_vector = std::vector>; +using Table = cudf::table; + +struct JoinTest : public cudf::test::BaseFixture { +}; + +std::pair, std::unique_ptr> get_saj_tables( + std::vector const& left_is_human_nulls, std::vector const& right_is_human_nulls) +{ + column_wrapper col0_0{{99, 1, 2, 0, 2}, {0, 1, 1, 1, 1}}; + strcol_wrapper col0_1({"s1", "s1", "s0", "s4", "s0"}, {1, 1, 0, 1, 1}); + column_wrapper col0_2{{0, 1, 2, 4, 1}}; + auto col0_names_col = strcol_wrapper{ + "Samuel Vimes", "Carrot Ironfoundersson", "Detritus", "Samuel Vimes", "Angua von Überwald"}; + auto col0_ages_col = column_wrapper{{48, 27, 351, 31, 25}}; + + auto col0_is_human_col = + column_wrapper{{true, true, false, false, false}, left_is_human_nulls.begin()}; + + auto col0_3 = cudf::test::structs_column_wrapper{ + {col0_names_col, col0_ages_col, col0_is_human_col}, {1, 1, 1, 1, 1}}; + + column_wrapper col1_0{{2, 2, 0, 4, -99}, {1, 1, 1, 1, 0}}; + strcol_wrapper col1_1({"s1", "s0", "s1", "s2", "s1"}); + column_wrapper col1_2{{1, 0, 1, 2, 1}, {1, 0, 1, 1, 1}}; + auto col1_names_col = strcol_wrapper{"Carrot Ironfoundersson", + "Angua von Überwald", + "Detritus", + "Carrot Ironfoundersson", + "Samuel Vimes"}; + auto col1_ages_col = column_wrapper{{351, 25, 27, 31, 48}}; + + auto col1_is_human_col = + column_wrapper{{true, false, false, false, true}, right_is_human_nulls.begin()}; + + auto col1_3 = + cudf::test::structs_column_wrapper{{col1_names_col, col1_ages_col, col1_is_human_col}}; + + column_vector cols0, cols1; + cols0.push_back(col0_0.release()); + cols0.push_back(col0_1.release()); + cols0.push_back(col0_2.release()); + cols0.push_back(col0_3.release()); + cols1.push_back(col1_0.release()); + cols1.push_back(col1_1.release()); + cols1.push_back(col1_2.release()); + cols1.push_back(col1_3.release()); + + return {std::make_unique(std::move(cols0)), std::make_unique
(std::move(cols1))}; +} + +TEST_F(JoinTest, SemiJoinWithStructsAndNulls) +{ + auto tables = get_saj_tables({1, 1, 0, 1, 0}, {1, 0, 0, 1, 1}); + + auto result = cudf::left_semi_join( + *tables.first, *tables.second, {0, 1, 3}, {0, 1, 3}, cudf::null_equality::EQUAL); + auto result_sort_order = cudf::sorted_order(result->view()); + auto sorted_result = cudf::gather(result->view(), *result_sort_order); + + column_wrapper col_gold_0{{99, 2}, {0, 1}}; + strcol_wrapper col_gold_1({"s1", "s0"}, {1, 1}); + column_wrapper col_gold_2{{0, 1}}; + auto col_gold_3_names_col = strcol_wrapper{"Samuel Vimes", "Angua von Überwald"}; + auto col_gold_3_ages_col = column_wrapper{{48, 25}}; + + auto col_gold_3_is_human_col = column_wrapper{{true, false}, {1, 0}}; + + auto col_gold_3 = cudf::test::structs_column_wrapper{ + {col_gold_3_names_col, col_gold_3_ages_col, col_gold_3_is_human_col}}; + + column_vector cols_gold; + cols_gold.push_back(col_gold_0.release()); + cols_gold.push_back(col_gold_1.release()); + cols_gold.push_back(col_gold_2.release()); + cols_gold.push_back(col_gold_3.release()); + Table gold(std::move(cols_gold)); + + auto gold_sort_order = cudf::sorted_order(gold.view()); + auto sorted_gold = cudf::gather(gold.view(), *gold_sort_order); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(*sorted_gold, *sorted_result); +} + +TEST_F(JoinTest, SemiJoinWithStructsAndNullsNotEqual) +{ + auto tables = get_saj_tables({1, 1, 0, 1, 1}, {1, 1, 0, 1, 1}); + + auto result = cudf::left_semi_join( + *tables.first, *tables.second, {0, 1, 3}, {0, 1, 3}, cudf::null_equality::UNEQUAL); + auto result_sort_order = cudf::sorted_order(result->view()); + auto sorted_result = cudf::gather(result->view(), *result_sort_order); + + column_wrapper col_gold_0{{2}, {1}}; + strcol_wrapper col_gold_1({"s0"}, {1}); + column_wrapper col_gold_2{{1}}; + auto col_gold_3_names_col = strcol_wrapper{"Angua von Überwald"}; + auto col_gold_3_ages_col = column_wrapper{{25}}; + + auto col_gold_3_is_human_col = column_wrapper{{false}, {1}}; + + auto col_gold_3 = cudf::test::structs_column_wrapper{ + {col_gold_3_names_col, col_gold_3_ages_col, col_gold_3_is_human_col}}; + + column_vector cols_gold; + cols_gold.push_back(col_gold_0.release()); + cols_gold.push_back(col_gold_1.release()); + cols_gold.push_back(col_gold_2.release()); + cols_gold.push_back(col_gold_3.release()); + Table gold(std::move(cols_gold)); + + auto gold_sort_order = cudf::sorted_order(gold.view()); + auto sorted_gold = cudf::gather(gold.view(), *gold_sort_order); + + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(*sorted_gold, *sorted_result); +} + +TEST_F(JoinTest, AntiJoinWithStructsAndNulls) +{ + auto tables = get_saj_tables({1, 1, 0, 1, 0}, {1, 0, 0, 1, 1}); + + auto result = cudf::left_anti_join( + *tables.first, *tables.second, {0, 1, 3}, {0, 1, 3}, cudf::null_equality::EQUAL); + auto result_sort_order = cudf::sorted_order(result->view()); + auto sorted_result = cudf::gather(result->view(), *result_sort_order); + + column_wrapper col_gold_0{{1, 2, 0}, {1, 1, 1}}; + strcol_wrapper col_gold_1({"s1", "s0", "s4"}, {1, 0, 1}); + column_wrapper col_gold_2{{1, 2, 4}}; + auto col_gold_3_names_col = strcol_wrapper{"Carrot Ironfoundersson", "Detritus", "Samuel Vimes"}; + auto col_gold_3_ages_col = column_wrapper{{27, 351, 31}}; + + auto col_gold_3_is_human_col = column_wrapper{{true, false, false}, {1, 0, 1}}; + + auto col_gold_3 = cudf::test::structs_column_wrapper{ + {col_gold_3_names_col, col_gold_3_ages_col, col_gold_3_is_human_col}}; + + column_vector cols_gold; + cols_gold.push_back(col_gold_0.release()); + cols_gold.push_back(col_gold_1.release()); + cols_gold.push_back(col_gold_2.release()); + cols_gold.push_back(col_gold_3.release()); + Table gold(std::move(cols_gold)); + + auto gold_sort_order = cudf::sorted_order(gold.view()); + auto sorted_gold = cudf::gather(gold.view(), *gold_sort_order); + + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(*sorted_gold, *sorted_result); +} + +TEST_F(JoinTest, AntiJoinWithStructsAndNullsNotEqual) +{ + auto tables = get_saj_tables({1, 1, 0, 1, 1}, {1, 1, 0, 1, 1}); + + auto result = cudf::left_anti_join( + *tables.first, *tables.second, {0, 1, 3}, {0, 1, 3}, cudf::null_equality::UNEQUAL); + auto result_sort_order = cudf::sorted_order(result->view()); + auto sorted_result = cudf::gather(result->view(), *result_sort_order); + + column_wrapper col_gold_0{{99, 1, 2, 0}, {0, 1, 1, 1}}; + strcol_wrapper col_gold_1({"s1", "s1", "s0", "s4"}, {1, 1, 0, 1}); + column_wrapper col_gold_2{{0, 1, 2, 4}}; + auto col_gold_3_names_col = + strcol_wrapper{"Samuel Vimes", "Carrot Ironfoundersson", "Detritus", "Samuel Vimes"}; + auto col_gold_3_ages_col = column_wrapper{{48, 27, 351, 31}}; + + auto col_gold_3_is_human_col = column_wrapper{{true, true, false, false}, {1, 1, 0, 1}}; + + auto col_gold_3 = cudf::test::structs_column_wrapper{ + {col_gold_3_names_col, col_gold_3_ages_col, col_gold_3_is_human_col}}; + + column_vector cols_gold; + cols_gold.push_back(col_gold_0.release()); + cols_gold.push_back(col_gold_1.release()); + cols_gold.push_back(col_gold_2.release()); + cols_gold.push_back(col_gold_3.release()); + Table gold(std::move(cols_gold)); + + auto gold_sort_order = cudf::sorted_order(gold.view()); + auto sorted_gold = cudf::gather(gold.view(), *gold_sort_order); + + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(*sorted_gold, *sorted_result); +} diff --git a/cpp/tests/join/semi_join_tests.cpp b/cpp/tests/join/semi_join_tests.cpp deleted file mode 100644 index 178a26dfdba..00000000000 --- a/cpp/tests/join/semi_join_tests.cpp +++ /dev/null @@ -1,109 +0,0 @@ -/* - * Copyright (c) 2019-2021, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include -#include -#include - -#include - -template -using column_wrapper = cudf::test::fixed_width_column_wrapper; -using strcol_wrapper = cudf::test::strings_column_wrapper; -using column_vector = std::vector>; -using Table = cudf::table; - -struct JoinTest : public cudf::test::BaseFixture { -}; - -TEST_F(JoinTest, SemiJoinWithStructsAndNulls) -{ - column_wrapper col0_0{{3, 1, 2, 0, 2}}; - strcol_wrapper col0_1({"s1", "s1", "s0", "s4", "s0"}, {1, 1, 0, 1, 1}); - column_wrapper col0_2{{0, 1, 2, 4, 1}}; - auto col0_names_col = strcol_wrapper{ - "Samuel Vimes", "Carrot Ironfoundersson", "Detritus", "Samuel Vimes", "Angua von Überwald"}; - auto col0_ages_col = column_wrapper{{48, 27, 351, 31, 25}}; - - auto col0_is_human_col = column_wrapper{{true, true, false, false, false}, {1, 1, 0, 1, 0}}; - - auto col0_3 = cudf::test::structs_column_wrapper{ - {col0_names_col, col0_ages_col, col0_is_human_col}, {1, 1, 1, 1, 1}}; - - column_wrapper col1_0{{2, 2, 0, 4, 3}}; - strcol_wrapper col1_1({"s1", "s0", "s1", "s2", "s1"}); - column_wrapper col1_2{{1, 0, 1, 2, 1}, {1, 0, 1, 1, 1}}; - auto col1_names_col = strcol_wrapper{"Carrot Ironfoundersson", - "Angua von Überwald", - "Detritus", - "Carrot Ironfoundersson", - "Samuel Vimes"}; - auto col1_ages_col = column_wrapper{{351, 25, 27, 31, 48}}; - - auto col1_is_human_col = column_wrapper{{true, false, false, false, true}, {1, 0, 0, 1, 1}}; - - auto col1_3 = - cudf::test::structs_column_wrapper{{col1_names_col, col1_ages_col, col1_is_human_col}}; - - column_vector cols0, cols1; - cols0.push_back(col0_0.release()); - cols0.push_back(col0_1.release()); - cols0.push_back(col0_2.release()); - cols0.push_back(col0_3.release()); - cols1.push_back(col1_0.release()); - cols1.push_back(col1_1.release()); - cols1.push_back(col1_2.release()); - cols1.push_back(col1_3.release()); - - Table t0(std::move(cols0)); - Table t1(std::move(cols1)); - - auto result = cudf::left_semi_join(t0, t1, {0, 1, 3}, {0, 1, 3}); - auto result_sort_order = cudf::sorted_order(result->view()); - auto sorted_result = cudf::gather(result->view(), *result_sort_order); - - column_wrapper col_gold_0{{3, 2}}; - strcol_wrapper col_gold_1({"s1", "s0"}, {1, 1}); - column_wrapper col_gold_2{{0, 1}}; - auto col_gold_3_names_col = strcol_wrapper{"Samuel Vimes", "Angua von Überwald"}; - auto col_gold_3_ages_col = column_wrapper{{48, 25}}; - - auto col_gold_3_is_human_col = column_wrapper{{true, false}, {1, 0}}; - - auto col_gold_3 = cudf::test::structs_column_wrapper{ - {col_gold_3_names_col, col_gold_3_ages_col, col_gold_3_is_human_col}}; - - column_vector cols_gold; - cols_gold.push_back(col_gold_0.release()); - cols_gold.push_back(col_gold_1.release()); - cols_gold.push_back(col_gold_2.release()); - cols_gold.push_back(col_gold_3.release()); - Table gold(std::move(cols_gold)); - - auto gold_sort_order = cudf::sorted_order(gold.view()); - auto sorted_gold = cudf::gather(gold.view(), *gold_sort_order); - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(*sorted_gold, *sorted_result); -} From 55289778ea17d638692934c40b2fc7bd332c04ab Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Thu, 20 May 2021 11:29:21 -0500 Subject: [PATCH 6/7] Small PR review tweaks. --- cpp/benchmarks/join/join_benchmark.cu | 9 ++++----- cpp/src/join/semi_join.cu | 4 ++-- 2 files changed, 6 insertions(+), 7 deletions(-) diff --git a/cpp/benchmarks/join/join_benchmark.cu b/cpp/benchmarks/join/join_benchmark.cu index 357852629f4..a7c109db9b4 100644 --- a/cpp/benchmarks/join/join_benchmark.cu +++ b/cpp/benchmarks/join/join_benchmark.cu @@ -38,7 +38,7 @@ class Join : public cudf::benchmark { }; template -static void BM_join(Join JoinFunc, benchmark::State& state) +static void BM_join(benchmark::State& state, Join JoinFunc) { const cudf::size_type build_table_size{(cudf::size_type)state.range(0)}; const cudf::size_type probe_table_size{(cudf::size_type)state.range(1)}; @@ -121,7 +121,7 @@ static void BM_join(Join JoinFunc, benchmark::State& state) cudf::null_equality compare_nulls) { \ return cudf::inner_join(left, right, left_on, right_on, compare_nulls); \ }; \ - BM_join(join, st); \ + BM_join(st, join); \ } JOIN_BENCHMARK_DEFINE(join_32bit, int32_t, int32_t, false); @@ -140,7 +140,7 @@ JOIN_BENCHMARK_DEFINE(join_64bit_nulls, int64_t, int64_t, true); cudf::null_equality compare_nulls) { \ return cudf::left_anti_join(left, right, left_on, right_on, compare_nulls); \ }; \ - BM_join(join, st); \ + BM_join(st, join); \ } LEFT_ANTI_JOIN_BENCHMARK_DEFINE(left_anti_join_32bit, int32_t, int32_t, false); @@ -159,7 +159,7 @@ LEFT_ANTI_JOIN_BENCHMARK_DEFINE(left_anti_join_64bit_nulls, int64_t, int64_t, tr cudf::null_equality compare_nulls) { \ return cudf::left_semi_join(left, right, left_on, right_on, compare_nulls); \ }; \ - BM_join(join, st); \ + BM_join(st, join); \ } LEFT_SEMI_JOIN_BENCHMARK_DEFINE(left_semi_join_32bit, int32_t, int32_t, false); @@ -277,4 +277,3 @@ BENCHMARK_REGISTER_F(Join, left_semi_join_64bit_nulls) ->Args({50'000'000, 50'000'000}) ->Args({40'000'000, 120'000'000}) ->UseManualTime(); - \ No newline at end of file diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 5fc3a3dbcda..725fff4cfdf 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -93,9 +93,9 @@ std::unique_ptr> left_semi_anti_join( auto hash_table = *hash_table_ptr; // if compare_nulls == NOT_EQUAL, we can simply ignore any rows that are - // entirely NULL as they will never compare to equal. + // contain a NULL in any column as they will never compare to equal. auto const row_bitmask = (compare_nulls == null_equality::EQUAL) - ? rmm::device_buffer{0, stream} + ? rmm::device_buffer{} : cudf::detail::bitmask_and(right_flattened_keys, stream); // skip rows that are null here. thrust::for_each_n( From 0c7377483d18e840a5668d484f22a8ab7ea4ae9d Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Thu, 20 May 2021 15:25:05 -0500 Subject: [PATCH 7/7] Comment grammar fix. --- cpp/src/join/semi_join.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 725fff4cfdf..cc34aed33ea 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -92,7 +92,7 @@ std::unique_ptr> left_semi_anti_join( equality_build); auto hash_table = *hash_table_ptr; - // if compare_nulls == NOT_EQUAL, we can simply ignore any rows that are + // if compare_nulls == UNEQUAL, we can simply ignore any rows that // contain a NULL in any column as they will never compare to equal. auto const row_bitmask = (compare_nulls == null_equality::EQUAL) ? rmm::device_buffer{}