From 1e1086e59e8a17fa66ffd60c22ed13d29e3eb8c1 Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Tue, 8 Dec 2020 19:11:57 +0000 Subject: [PATCH 01/14] Initial pass at fixing nulls not equal performance with joins --- cpp/benchmarks/join/join_benchmark.cu | 88 +++++++++++++++++++++++---- cpp/include/cudf/join.hpp | 2 + cpp/src/join/hash_join.cu | 24 ++++++-- cpp/src/join/hash_join.cuh | 2 + cpp/src/join/join.cu | 11 ++-- cpp/src/join/join_kernels.cuh | 28 +++++---- cpp/tests/join/join_tests.cpp | 2 +- 7 files changed, 124 insertions(+), 33 deletions(-) diff --git a/cpp/benchmarks/join/join_benchmark.cu b/cpp/benchmarks/join/join_benchmark.cu index b18ceafdae6..3ca0a413665 100644 --- a/cpp/benchmarks/join/join_benchmark.cu +++ b/cpp/benchmarks/join/join_benchmark.cu @@ -36,7 +36,7 @@ template class Join : public cudf::benchmark { }; -template +template static void BM_join(benchmark::State &state) { const cudf::size_type build_table_size{(cudf::size_type)state.range(0)}; @@ -47,10 +47,31 @@ static void BM_join(benchmark::State &state) // Generate build and probe tables - auto build_key_column = - cudf::make_numeric_column(cudf::data_type(cudf::type_to_id()), build_table_size); - auto probe_key_column = - cudf::make_numeric_column(cudf::data_type(cudf::type_to_id()), probe_table_size); + auto build_random_null_mask = [](int size) { + if (Nullable) { + // roughly 25% nulls + auto validity = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), [](auto i) { return (std::rand() & 3) == 0; }); + return cudf::test::detail::make_null_mask(validity, validity + size); + } else { + return cudf::create_null_mask(size, cudf::mask_state::UNINITIALIZED); + } + }; + + std::unique_ptr build_key_column = [&]() { + return Nullable ? cudf::make_numeric_column(cudf::data_type(cudf::type_to_id()), + build_table_size, + build_random_null_mask(build_table_size)) + : cudf::make_numeric_column(cudf::data_type(cudf::type_to_id()), + build_table_size); + }(); + std::unique_ptr probe_key_column = [&]() { + return Nullable ? cudf::make_numeric_column(cudf::data_type(cudf::type_to_id()), + probe_table_size, + build_random_null_mask(probe_table_size)) + : cudf::make_numeric_column(cudf::data_type(cudf::type_to_id()), + probe_table_size); + }(); generate_input_tables( build_key_column->mutable_view().data(), @@ -73,6 +94,25 @@ static void BM_join(benchmark::State &state) cudf::table_view build_table({build_key_column->view(), build_payload_column}); cudf::table_view probe_table({probe_key_column->view(), probe_payload_column}); + /* + auto build_column = [](int table_size, int null_count) { + auto data_it = thrust::make_counting_iterator(0); + if (Nullable) { + auto validity = thrust::make_transform_iterator(data_it, [](auto i) { return i % 3 == 0; }); + + return cudf::test::fixed_width_column_wrapper( + data_it, data_it + table_size, validity); + } else { + return cudf::test::fixed_width_column_wrapper(data_it, data_it + table_size); + } + }; + + CHECK_CUDA(0); + + cudf::table_view build_table({build_key_column->view(), build_column(build_table_size, 3)}); + cudf::table_view probe_table({probe_key_column->view(), build_column(probe_table_size, 5)}); +*/ + // Setup join parameters and result table std::vector columns_to_join = {0}; @@ -82,17 +122,23 @@ static void BM_join(benchmark::State &state) for (auto _ : state) { cuda_event_timer raii(state, true, 0); - auto result = - cudf::inner_join(probe_table, build_table, columns_to_join, columns_to_join, {{0, 0}}); + auto result = cudf::inner_join(probe_table, + build_table, + columns_to_join, + columns_to_join, + {{0, 0}}, + cudf::null_equality::UNEQUAL); } } -#define JOIN_BENCHMARK_DEFINE(name, key_type, payload_type) \ - 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) { BM_join(st); } -JOIN_BENCHMARK_DEFINE(join_32bit, int32_t, int32_t); -JOIN_BENCHMARK_DEFINE(join_64bit, int64_t, int64_t); +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); BENCHMARK_REGISTER_F(Join, join_32bit) ->Unit(benchmark::kMillisecond) @@ -111,3 +157,21 @@ BENCHMARK_REGISTER_F(Join, join_64bit) ->Args({50'000'000, 50'000'000}) ->Args({40'000'000, 120'000'000}) ->UseManualTime(); + +BENCHMARK_REGISTER_F(Join, 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'000, 100'000'000}) + ->Args({100'000'000, 100'000'000}) + ->Args({80'000'000, 240'000'000}) + ->UseManualTime(); + +BENCHMARK_REGISTER_F(Join, join_64bit_nulls) + ->Unit(benchmark::kMillisecond) + ->Args({50'000'000, 50'000'000}) + ->Args({40'000'000, 120'000'000}) + ->UseManualTime(); diff --git a/cpp/include/cudf/join.hpp b/cpp/include/cudf/join.hpp index 37847c41339..b2c1296ccef 100644 --- a/cpp/include/cudf/join.hpp +++ b/cpp/include/cudf/join.hpp @@ -394,10 +394,12 @@ class hash_join { * * @param build The build table, from which the hash table is built. * @param build_on The column indices from `build` to join on. + * @param compare_nulls Controls whether null join-key values should match or not. * @param stream CUDA stream used for device memory operations and kernel launches */ hash_join(cudf::table_view const& build, std::vector const& build_on, + null_equality compare_nulls, rmm::cuda_stream_view stream = rmm::cuda_stream_default); /** diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 67b9d3436d8..17989031cd3 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -196,13 +196,18 @@ get_left_join_indices_complement(rmm::device_vector &right_indices, * @throw std::out_of_range if elements of `build_on` exceed the number of columns in the `build` * table. * - * @param build_table Table of build side columns to join. + * @param build_table Device view of table of build side columns to join. + * @param build Table of build side columns. + * @param compare_nulls Controls whether null join-key values should match or not. * @param stream CUDA stream used for device memory operations and kernel launches. * * @return Built hash table. */ std::unique_ptr> build_join_hash_table( - cudf::table_device_view build_table, rmm::cuda_stream_view stream) + cudf::table_device_view build_table, + cudf::table_view const &build, + null_equality compare_nulls, + rmm::cuda_stream_view stream) { CUDF_EXPECTS(0 != build_table.num_columns(), "Selected build dataset is empty"); CUDF_EXPECTS(0 != build_table.num_rows(), "Build side table has no rows"); @@ -221,8 +226,18 @@ std::unique_ptr> build_join_ rmm::device_scalar failure(0, stream); constexpr int block_size{DEFAULT_JOIN_BLOCK_SIZE}; detail::grid_1d config(build_table_num_rows, block_size); + auto row_bitmask = [compare_nulls, &stream, &build]() { + if (compare_nulls == null_equality::EQUAL) + return rmm::device_buffer{0, stream}; + else + return cudf::detail::bitmask_and(build, stream); + }(); build_hash_table<<>>( - *hash_table, hash_build, build_table_num_rows, failure.data()); + *hash_table, + hash_build, + build_table_num_rows, + static_cast(row_bitmask.data()), + failure.data()); // Check error code from the kernel if (failure.value(stream) == 1) { CUDF_FAIL("Hash Table insert failure."); } @@ -484,6 +499,7 @@ hash_join::hash_join_impl::~hash_join_impl() = default; hash_join::hash_join_impl::hash_join_impl(cudf::table_view const &build, std::vector const &build_on, + null_equality compare_nulls, rmm::cuda_stream_view stream) : _build(build), _build_selected(build.select(build_on)), @@ -498,7 +514,7 @@ hash_join::hash_join_impl::hash_join_impl(cudf::table_view const &build, if (_build_on.empty() || 0 == build.num_rows()) { return; } auto build_table = cudf::table_device_view::create(_build_selected, stream); - _hash_table = build_join_hash_table(*build_table, stream); + _hash_table = build_join_hash_table(*build_table, build, compare_nulls, stream); } std::pair, std::unique_ptr> diff --git a/cpp/src/join/hash_join.cuh b/cpp/src/join/hash_join.cuh index 36cb486d4c1..b9f73a8214e 100644 --- a/cpp/src/join/hash_join.cuh +++ b/cpp/src/join/hash_join.cuh @@ -226,9 +226,11 @@ struct hash_join::hash_join_impl { * * @param build The build table, from which the hash table is built. * @param build_on The column indices from `build` to join on. + * @param compare_nulls Controls whether null join-key values should match or not. */ hash_join_impl(cudf::table_view const& build, std::vector const& build_on, + null_equality compare_nulls, rmm::cuda_stream_view stream = rmm::cuda_stream_default); std::pair, std::unique_ptr> inner_join( diff --git a/cpp/src/join/join.cu b/cpp/src/join/join.cu index af649fe5fb0..ce27cfcd616 100644 --- a/cpp/src/join/join.cu +++ b/cpp/src/join/join.cu @@ -51,7 +51,7 @@ std::unique_ptr inner_join( // building/probing the hash map. Because building is typically more expensive than probing, we // build the hash map from the smaller table. if (right.num_rows() > left.num_rows()) { - cudf::hash_join hj_obj(left, left_on, stream); + cudf::hash_join hj_obj(left, left_on, compare_nulls, stream); auto actual_columns_in_common = columns_in_common; std::for_each(actual_columns_in_common.begin(), actual_columns_in_common.end(), [](auto& pair) { std::swap(pair.first, pair.second); @@ -66,7 +66,7 @@ std::unique_ptr
inner_join( return cudf::detail::combine_table_pair(std::move(probe_build_pair.second), std::move(probe_build_pair.first)); } else { - cudf::hash_join hj_obj(right, right_on, stream); + cudf::hash_join hj_obj(right, right_on, compare_nulls, stream); auto probe_build_pair = hj_obj.inner_join(left, left_on, columns_in_common, @@ -99,7 +99,7 @@ std::unique_ptr
left_join( table_view const left = scatter_columns(matched.second.front(), left_on, left_input); table_view const right = scatter_columns(matched.second.back(), right_on, right_input); - cudf::hash_join hj_obj(right, right_on, stream); + cudf::hash_join hj_obj(right, right_on, compare_nulls, stream); return hj_obj.left_join(left, left_on, columns_in_common, compare_nulls, stream, mr); } @@ -123,7 +123,7 @@ std::unique_ptr
full_join( table_view const left = scatter_columns(matched.second.front(), left_on, left_input); table_view const right = scatter_columns(matched.second.back(), right_on, right_input); - cudf::hash_join hj_obj(right, right_on, stream); + cudf::hash_join hj_obj(right, right_on, compare_nulls, stream); return hj_obj.full_join(left, left_on, columns_in_common, compare_nulls, stream, mr); } @@ -133,8 +133,9 @@ hash_join::~hash_join() = default; hash_join::hash_join(cudf::table_view const& build, std::vector const& build_on, + null_equality compare_nulls, rmm::cuda_stream_view stream) - : impl{std::make_unique(build, build_on, stream)} + : impl{std::make_unique(build, build_on, compare_nulls, stream)} { } diff --git a/cpp/src/join/join_kernels.cuh b/cpp/src/join/join_kernels.cuh index 02e7b6c6c85..130e7e053d6 100644 --- a/cpp/src/join/join_kernels.cuh +++ b/cpp/src/join/join_kernels.cuh @@ -17,6 +17,7 @@ #pragma once #include +#include #include #include @@ -71,28 +72,33 @@ constexpr auto remap_sentinel_hash(H hash, S sentinel) * @param[in,out] multi_map The hash table to be built to insert rows into * @param[in] hash_build Row hasher for the build table * @param[in] build_table_num_rows The number of rows in the build table + * @param[in] row_bitmask Bitmask where bit `i` indicates the presence of a null + * value in row `i` of input keys. This is nullptr if nulls are equal. * @param[out] error Pointer used to set an error code if the insert fails */ template __global__ void build_hash_table(multimap_type multi_map, row_hash hash_build, const cudf::size_type build_table_num_rows, + bitmask_type const* row_bitmask, int* error) { cudf::size_type i = threadIdx.x + blockIdx.x * blockDim.x; while (i < build_table_num_rows) { - // Compute the hash value of this row - auto const row_hash_value = remap_sentinel_hash(hash_build(i), multi_map.get_unused_key()); - - // Insert the (row hash value, row index) into the map - // using the row hash value to determine the location in the - // hash map where the new pair should be inserted - const auto insert_location = - multi_map.insert(thrust::make_pair(row_hash_value, i), true, row_hash_value); - - // If the insert failed, set the error code accordingly - if (multi_map.end() == insert_location) { *error = 1; } + if (!row_bitmask || cudf::bit_is_set(row_bitmask, i)) { + // Compute the hash value of this row + auto const row_hash_value = remap_sentinel_hash(hash_build(i), multi_map.get_unused_key()); + + // Insert the (row hash value, row index) into the map + // using the row hash value to determine the location in the + // hash map where the new pair should be inserted + const auto insert_location = + multi_map.insert(thrust::make_pair(row_hash_value, i), true, row_hash_value); + + // If the insert failed, set the error code accordingly + if (multi_map.end() == insert_location) { *error = 1; } + } i += blockDim.x * gridDim.x; } } diff --git a/cpp/tests/join/join_tests.cpp b/cpp/tests/join/join_tests.cpp index 406bc69220a..efc5330ea7d 100644 --- a/cpp/tests/join/join_tests.cpp +++ b/cpp/tests/join/join_tests.cpp @@ -1120,7 +1120,7 @@ TEST_F(JoinTest, HashJoinSequentialProbes) Table t1(std::move(cols1)); - cudf::hash_join hash_join(t1, {0, 1}); + cudf::hash_join hash_join(t1, {0, 1}, cudf::null_equality::EQUAL); { CVector cols0; From 95f791f3d7d6557126bceee3dcf1bb61e1703aae Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Tue, 8 Dec 2020 19:38:26 +0000 Subject: [PATCH 02/14] adding changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 52210a69ef0..90ff975f0e3 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -14,6 +14,7 @@ ## Bug Fixes - PR #6912 Fix rmm_mode=managed parameter for gtests +- PR #6943 Fix join with nulls not equal performance # cuDF 0.17.0 (Date TBD) From 7db4f1b0533c8116f0f8094656f2ffef7110e0d3 Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Tue, 8 Dec 2020 17:28:13 -0500 Subject: [PATCH 03/14] Apply suggestions from code review Thanks for the suggestions! Co-authored-by: Jake Hemstad --- cpp/src/join/hash_join.cu | 2 +- cpp/src/join/join_kernels.cuh | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 000a4b8c441..b66623c5dc5 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -226,7 +226,7 @@ std::unique_ptr> build_join_ rmm::device_scalar failure(0, stream); constexpr int block_size{DEFAULT_JOIN_BLOCK_SIZE}; detail::grid_1d config(build_table_num_rows, block_size); - auto row_bitmask = [compare_nulls, &stream, &build]() { + auto const row_bitmask = [compare_nulls, &stream, &build]() { if (compare_nulls == null_equality::EQUAL) return rmm::device_buffer{0, stream}; else diff --git a/cpp/src/join/join_kernels.cuh b/cpp/src/join/join_kernels.cuh index 130e7e053d6..c353ec2e895 100644 --- a/cpp/src/join/join_kernels.cuh +++ b/cpp/src/join/join_kernels.cuh @@ -93,7 +93,7 @@ __global__ void build_hash_table(multimap_type multi_map, // Insert the (row hash value, row index) into the map // using the row hash value to determine the location in the // hash map where the new pair should be inserted - const auto insert_location = + auto const insert_location = multi_map.insert(thrust::make_pair(row_hash_value, i), true, row_hash_value); // If the insert failed, set the error code accordingly From 539f46150323bfc3398b03fa98962eedee6d6f8c Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Wed, 9 Dec 2020 05:57:26 +0000 Subject: [PATCH 04/14] removing stale comments --- cpp/benchmarks/join/join_benchmark.cu | 19 ------------------- 1 file changed, 19 deletions(-) diff --git a/cpp/benchmarks/join/join_benchmark.cu b/cpp/benchmarks/join/join_benchmark.cu index 3ca0a413665..10991589723 100644 --- a/cpp/benchmarks/join/join_benchmark.cu +++ b/cpp/benchmarks/join/join_benchmark.cu @@ -94,25 +94,6 @@ static void BM_join(benchmark::State &state) cudf::table_view build_table({build_key_column->view(), build_payload_column}); cudf::table_view probe_table({probe_key_column->view(), probe_payload_column}); - /* - auto build_column = [](int table_size, int null_count) { - auto data_it = thrust::make_counting_iterator(0); - if (Nullable) { - auto validity = thrust::make_transform_iterator(data_it, [](auto i) { return i % 3 == 0; }); - - return cudf::test::fixed_width_column_wrapper( - data_it, data_it + table_size, validity); - } else { - return cudf::test::fixed_width_column_wrapper(data_it, data_it + table_size); - } - }; - - CHECK_CUDA(0); - - cudf::table_view build_table({build_key_column->view(), build_column(build_table_size, 3)}); - cudf::table_view probe_table({probe_key_column->view(), build_column(probe_table_size, 5)}); -*/ - // Setup join parameters and result table std::vector columns_to_join = {0}; From 8af46b2851bdcdfd2da255ebe5b872dbffcaa574 Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Wed, 9 Dec 2020 06:03:15 +0000 Subject: [PATCH 05/14] attempting to clarify parameter comments --- cpp/src/join/hash_join.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index b66623c5dc5..802907828b5 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -196,8 +196,8 @@ get_left_join_indices_complement(rmm::device_vector &right_indices, * @throw std::out_of_range if elements of `build_on` exceed the number of columns in the `build` * table. * - * @param build_table Device view of table of build side columns to join. - * @param build Table of build side columns. + * @param build_table Device view of table of columns used to build join hash. + * @param build Table of columns used to build join hash. * @param compare_nulls Controls whether null join-key values should match or not. * @param stream CUDA stream used for device memory operations and kernel launches. * From e1ee434dc542dcb827fd228322dafe7779c89d1a Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Thu, 10 Dec 2020 02:10:07 +0000 Subject: [PATCH 06/14] moving device view creation up into build_join_hash_table and only passing one table --- cpp/src/join/hash_join.cu | 17 ++++++++--------- 1 file changed, 8 insertions(+), 9 deletions(-) diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 802907828b5..febfd371e4c 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -196,7 +196,6 @@ get_left_join_indices_complement(rmm::device_vector &right_indices, * @throw std::out_of_range if elements of `build_on` exceed the number of columns in the `build` * table. * - * @param build_table Device view of table of columns used to build join hash. * @param build Table of columns used to build join hash. * @param compare_nulls Controls whether null join-key values should match or not. * @param stream CUDA stream used for device memory operations and kernel launches. @@ -204,15 +203,16 @@ get_left_join_indices_complement(rmm::device_vector &right_indices, * @return Built hash table. */ std::unique_ptr> build_join_hash_table( - cudf::table_device_view build_table, cudf::table_view const &build, null_equality compare_nulls, rmm::cuda_stream_view stream) { - CUDF_EXPECTS(0 != build_table.num_columns(), "Selected build dataset is empty"); - CUDF_EXPECTS(0 != build_table.num_rows(), "Build side table has no rows"); + auto build_table = cudf::table_device_view::create(build, stream); + + CUDF_EXPECTS(0 != build_table->num_columns(), "Selected build dataset is empty"); + CUDF_EXPECTS(0 != build_table->num_rows(), "Build side table has no rows"); - const size_type build_table_num_rows{build_table.num_rows()}; + const size_type build_table_num_rows{build_table->num_rows()}; size_t const hash_table_size = compute_hash_table_size(build_table_num_rows); auto hash_table = multimap_type::create(hash_table_size, @@ -222,7 +222,7 @@ std::unique_ptr> build_join_ multimap_type::key_equal(), multimap_type::allocator_type()); - row_hash hash_build{build_table}; + row_hash hash_build{*build_table}; rmm::device_scalar failure(0, stream); constexpr int block_size{DEFAULT_JOIN_BLOCK_SIZE}; detail::grid_1d config(build_table_num_rows, block_size); @@ -236,7 +236,7 @@ std::unique_ptr> build_join_ *hash_table, hash_build, build_table_num_rows, - static_cast(row_bitmask.data()), + static_cast(row_bitmask.data()), failure.data()); // Check error code from the kernel if (failure.value(stream) == 1) { CUDF_FAIL("Hash Table insert failure."); } @@ -515,8 +515,7 @@ hash_join::hash_join_impl::hash_join_impl(cudf::table_view const &build, if (_build_on.empty() || 0 == build.num_rows()) { return; } - auto build_table = cudf::table_device_view::create(_build_selected, stream); - _hash_table = build_join_hash_table(*build_table, build, compare_nulls, stream); + _hash_table = build_join_hash_table(_build_selected, compare_nulls, stream); } std::pair, std::unique_ptr> From 3d77c7328d6f30384ed45b4673ca4f1fe05c430e Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Thu, 10 Dec 2020 07:40:35 +0000 Subject: [PATCH 07/14] updating based on review --- cpp/src/join/hash_join.cu | 16 +++++++--------- 1 file changed, 7 insertions(+), 9 deletions(-) diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index febfd371e4c..61e558d06cf 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -203,16 +203,14 @@ get_left_join_indices_complement(rmm::device_vector &right_indices, * @return Built hash table. */ std::unique_ptr> build_join_hash_table( - cudf::table_view const &build, - null_equality compare_nulls, - rmm::cuda_stream_view stream) + cudf::table_view const &build, null_equality compare_nulls, rmm::cuda_stream_view stream) { - auto build_table = cudf::table_device_view::create(build, stream); + auto build_device_table = cudf::table_device_view::create(build, stream); - CUDF_EXPECTS(0 != build_table->num_columns(), "Selected build dataset is empty"); - CUDF_EXPECTS(0 != build_table->num_rows(), "Build side table has no rows"); + CUDF_EXPECTS(0 != build_device_table->num_columns(), "Selected build dataset is empty"); + CUDF_EXPECTS(0 != build_device_table->num_rows(), "Build side table has no rows"); - const size_type build_table_num_rows{build_table->num_rows()}; + const size_type build_table_num_rows{build_device_table->num_rows()}; size_t const hash_table_size = compute_hash_table_size(build_table_num_rows); auto hash_table = multimap_type::create(hash_table_size, @@ -222,7 +220,7 @@ std::unique_ptr> build_join_ multimap_type::key_equal(), multimap_type::allocator_type()); - row_hash hash_build{*build_table}; + row_hash hash_build{*build_device_table}; rmm::device_scalar failure(0, stream); constexpr int block_size{DEFAULT_JOIN_BLOCK_SIZE}; detail::grid_1d config(build_table_num_rows, block_size); @@ -515,7 +513,7 @@ hash_join::hash_join_impl::hash_join_impl(cudf::table_view const &build, if (_build_on.empty() || 0 == build.num_rows()) { return; } - _hash_table = build_join_hash_table(_build_selected, compare_nulls, stream); + _hash_table = build_join_hash_table(_build_selected, compare_nulls, stream); } std::pair, std::unique_ptr> From e2f28101e0e397a009747a5ee432fa574f3829f9 Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Thu, 10 Dec 2020 13:03:02 -0500 Subject: [PATCH 08/14] Apply suggestions from code review Thanks for the suggestions. Co-authored-by: nvdbaranec <56695930+nvdbaranec@users.noreply.github.com> --- cpp/src/join/hash_join.cu | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 61e558d06cf..1be554d0876 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -210,7 +210,7 @@ std::unique_ptr> build_join_ CUDF_EXPECTS(0 != build_device_table->num_columns(), "Selected build dataset is empty"); CUDF_EXPECTS(0 != build_device_table->num_rows(), "Build side table has no rows"); - const size_type build_table_num_rows{build_device_table->num_rows()}; + size_type const build_table_num_rows{build_device_table->num_rows()}; size_t const hash_table_size = compute_hash_table_size(build_table_num_rows); auto hash_table = multimap_type::create(hash_table_size, @@ -225,10 +225,7 @@ std::unique_ptr> build_join_ constexpr int block_size{DEFAULT_JOIN_BLOCK_SIZE}; detail::grid_1d config(build_table_num_rows, block_size); auto const row_bitmask = [compare_nulls, &stream, &build]() { - if (compare_nulls == null_equality::EQUAL) - return rmm::device_buffer{0, stream}; - else - return cudf::detail::bitmask_and(build, stream); + return compare_nulls == null_equality::EQUAL ? rmm::device_buffer{0, stream} : cudf::detail::bitmask_and(build, stream); }(); build_hash_table<<>>( *hash_table, From 9ab1e194e07620ac50dbabdcf007c4f98bf2a9ff Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Thu, 10 Dec 2020 18:29:54 +0000 Subject: [PATCH 09/14] Adding srand per review comments --- cpp/benchmarks/join/join_benchmark.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/benchmarks/join/join_benchmark.cu b/cpp/benchmarks/join/join_benchmark.cu index 10991589723..acd321a6ae6 100644 --- a/cpp/benchmarks/join/join_benchmark.cu +++ b/cpp/benchmarks/join/join_benchmark.cu @@ -46,7 +46,7 @@ static void BM_join(benchmark::State &state) const bool is_build_table_key_unique = true; // Generate build and probe tables - + std::srand(1337); auto build_random_null_mask = [](int size) { if (Nullable) { // roughly 25% nulls From 6b0e5d496eceef708fb8826e43ca35cc14be846e Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Thu, 10 Dec 2020 18:37:13 +0000 Subject: [PATCH 10/14] linting --- cpp/src/join/hash_join.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 1be554d0876..454a5a3c3ea 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -225,7 +225,8 @@ std::unique_ptr> build_join_ constexpr int block_size{DEFAULT_JOIN_BLOCK_SIZE}; detail::grid_1d config(build_table_num_rows, block_size); auto const row_bitmask = [compare_nulls, &stream, &build]() { - return compare_nulls == null_equality::EQUAL ? rmm::device_buffer{0, stream} : cudf::detail::bitmask_and(build, stream); + return compare_nulls == null_equality::EQUAL ? rmm::device_buffer{0, stream} + : cudf::detail::bitmask_and(build, stream); }(); build_hash_table<<>>( *hash_table, From 2e233da36056d62bcbfc2fdf8a27d78eb29d1aa7 Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Thu, 10 Dec 2020 19:36:40 +0000 Subject: [PATCH 11/14] review comment changes --- cpp/src/join/hash_join.cu | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 454a5a3c3ea..bbb20a103c4 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -224,10 +224,9 @@ std::unique_ptr> build_join_ rmm::device_scalar failure(0, stream); constexpr int block_size{DEFAULT_JOIN_BLOCK_SIZE}; detail::grid_1d config(build_table_num_rows, block_size); - auto const row_bitmask = [compare_nulls, &stream, &build]() { - return compare_nulls == null_equality::EQUAL ? rmm::device_buffer{0, stream} - : cudf::detail::bitmask_and(build, stream); - }(); + auto const row_bitmask = (compare_nulls == null_equality::EQUAL) + ? rmm::device_buffer{0, stream} + : cudf::detail::bitmask_and(build, stream); build_hash_table<<>>( *hash_table, hash_build, From e57855ad67db58dc12b7831fc4b17e112be8ed9f Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Fri, 11 Dec 2020 00:29:39 -0500 Subject: [PATCH 12/14] Apply suggestions from code review Co-authored-by: Mark Harris --- cpp/benchmarks/join/join_benchmark.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/benchmarks/join/join_benchmark.cu b/cpp/benchmarks/join/join_benchmark.cu index acd321a6ae6..47e74c7c9ff 100644 --- a/cpp/benchmarks/join/join_benchmark.cu +++ b/cpp/benchmarks/join/join_benchmark.cu @@ -46,12 +46,12 @@ static void BM_join(benchmark::State &state) const bool is_build_table_key_unique = true; // Generate build and probe tables - std::srand(1337); + cudf::test::UniformRandomGenerator rand_gen(0, build_table_size); auto build_random_null_mask = [](int size) { if (Nullable) { // roughly 25% nulls auto validity = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), [](auto i) { return (std::rand() & 3) == 0; }); + thrust::make_counting_iterator(0), [](auto i) { return (rand_gen.generate() & 3) == 0; }); return cudf::test::detail::make_null_mask(validity, validity + size); } else { return cudf::create_null_mask(size, cudf::mask_state::UNINITIALIZED); From a458a843539b9fbaa7caa638881fa29578ca08fe Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Fri, 11 Dec 2020 05:58:56 +0000 Subject: [PATCH 13/14] cleanup and fixes for UniformRandomGenerator --- cpp/benchmarks/join/join_benchmark.cu | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/cpp/benchmarks/join/join_benchmark.cu b/cpp/benchmarks/join/join_benchmark.cu index 47e74c7c9ff..ea97a47aa90 100644 --- a/cpp/benchmarks/join/join_benchmark.cu +++ b/cpp/benchmarks/join/join_benchmark.cu @@ -23,6 +23,7 @@ #include #include #include +#include #include #include @@ -46,12 +47,12 @@ static void BM_join(benchmark::State &state) const bool is_build_table_key_unique = true; // Generate build and probe tables - cudf::test::UniformRandomGenerator rand_gen(0, build_table_size); - auto build_random_null_mask = [](int size) { + cudf::test::UniformRandomGenerator rand_gen(0, build_table_size); + auto build_random_null_mask = [&rand_gen](int size) { if (Nullable) { // roughly 25% nulls auto validity = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), [](auto i) { return (rand_gen.generate() & 3) == 0; }); + thrust::make_counting_iterator(0), [&rand_gen](auto i) { return (rand_gen.generate() & 3) == 0; }); return cudf::test::detail::make_null_mask(validity, validity + size); } else { return cudf::create_null_mask(size, cudf::mask_state::UNINITIALIZED); From 39403b2129b9b35c12e9ab1c33f49f577f8442ed Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Fri, 11 Dec 2020 06:28:09 +0000 Subject: [PATCH 14/14] linting --- 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 ea97a47aa90..bd013afc451 100644 --- a/cpp/benchmarks/join/join_benchmark.cu +++ b/cpp/benchmarks/join/join_benchmark.cu @@ -52,7 +52,8 @@ static void BM_join(benchmark::State &state) if (Nullable) { // roughly 25% nulls auto validity = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), [&rand_gen](auto i) { return (rand_gen.generate() & 3) == 0; }); + thrust::make_counting_iterator(0), + [&rand_gen](auto i) { return (rand_gen.generate() & 3) == 0; }); return cudf::test::detail::make_null_mask(validity, validity + size); } else { return cudf::create_null_mask(size, cudf::mask_state::UNINITIALIZED);