From 4cc83025ddabb98933aeb9c7f9f657d1a0c96941 Mon Sep 17 00:00:00 2001 From: davidwendt Date: Thu, 29 Jul 2021 19:12:15 -0400 Subject: [PATCH 1/4] Move template param to function param in cudf::detail::left_semi_anti_join --- cpp/src/join/semi_join.cu | 88 ++++++++++++++++++++++----------------- 1 file changed, 49 insertions(+), 39 deletions(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index cc34aed33ea..788b489d4ca 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -18,15 +18,11 @@ #include #include -#include - #include #include #include -#include #include #include -#include #include #include @@ -34,11 +30,15 @@ #include #include +#include +#include +#include + namespace cudf { namespace detail { -template std::unique_ptr> left_semi_anti_join( + join_kind const kind, cudf::table_view const& left_keys, cudf::table_view const& right_keys, null_equality compare_nulls, @@ -48,13 +48,13 @@ std::unique_ptr> left_semi_anti_join( CUDF_EXPECTS(0 != left_keys.num_columns(), "Left table is empty"); CUDF_EXPECTS(0 != right_keys.num_columns(), "Right table is empty"); - if (is_trivial_join(left_keys, right_keys, JoinKind)) { + if (is_trivial_join(left_keys, right_keys, kind)) { return std::make_unique>(0, stream, mr); } - if ((join_kind::LEFT_ANTI_JOIN == JoinKind) && (0 == right_keys.num_rows())) { + if ((join_kind::LEFT_ANTI_JOIN == kind) && (0 == right_keys.num_rows())) { auto result = std::make_unique>(left_keys.num_rows(), stream, mr); - thrust::sequence(thrust::cuda::par.on(stream.value()), result->begin(), result->end()); + thrust::sequence(rmm::exec_policy(stream), result->begin(), result->end()); return result; } @@ -115,7 +115,7 @@ std::unique_ptr> left_semi_anti_join( // // For semi join we want contains to be true, for anti join we want contains to be false - bool join_type_boolean = (JoinKind == join_kind::LEFT_SEMI_JOIN); + bool const join_type_boolean = (kind == join_kind::LEFT_SEMI_JOIN); auto gather_map = std::make_unique>(left_num_rows, stream, mr); @@ -152,27 +152,26 @@ std::unique_ptr> left_semi_anti_join( * @throws cudf::logic_error if number of returned columns is 0 * @throws cudf::logic_error if number of elements in `right_on` and `left_on` are not equal * - * @param[in] left The left table - * @param[in] right The right table - * @param[in] left_on The column indices from `left` to join on. - * The column from `left` indicated by `left_on[i]` - * will be compared against the column from `right` - * indicated by `right_on[i]`. - * @param[in] right_on The column indices from `right` to join on. - * The column from `right` indicated by `right_on[i]` - * will be compared against the column from `left` - * indicated by `left_on[i]`. - * @param[in] compare_nulls Controls whether null join-key values should match or not. - * @param[in] mr Device memory resource to used to allocate the returned table's - * device memory - * @param[in] stream CUDA stream used for device memory operations and kernel launches. - * @tparam join_kind Indicates whether to do LEFT_SEMI_JOIN or LEFT_ANTI_JOIN + * @param kind Indicates whether to do LEFT_SEMI_JOIN or LEFT_ANTI_JOIN + * @param left The left table + * @param right The right table + * @param left_on The column indices from `left` to join on. + * The column from `left` indicated by `left_on[i]` + * will be compared against the column from `right` + * indicated by `right_on[i]`. + * @param right_on The column indices from `right` to join on. + * The column from `right` indicated by `right_on[i]` + * will be compared against the column from `left` + * indicated by `left_on[i]`. + * @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. + * @param mr Device memory resource to used to allocate the returned table * - * @returns Result of joining `left` and `right` tables on the columns - * specified by `left_on` and `right_on`. + * @returns Result of joining `left` and `right` tables on the columns + * specified by `left_on` and `right_on`. */ -template std::unique_ptr left_semi_anti_join( + join_kind const kind, cudf::table_view const& left, cudf::table_view const& right, std::vector const& left_on, @@ -183,11 +182,11 @@ std::unique_ptr left_semi_anti_join( { CUDF_EXPECTS(left_on.size() == right_on.size(), "Mismatch in number of columns to be joined on"); - if ((left_on.empty() || right_on.empty()) || is_trivial_join(left, right, JoinKind)) { + if ((left_on.empty() || right_on.empty()) || is_trivial_join(left, right, kind)) { return empty_like(left); } - if ((join_kind::LEFT_ANTI_JOIN == JoinKind) && (0 == right.num_rows())) { + if ((join_kind::LEFT_ANTI_JOIN == kind) && (0 == right.num_rows())) { // Everything matches, just copy the proper columns from the left table return std::make_unique(left, stream, mr); } @@ -202,8 +201,7 @@ std::unique_ptr left_semi_anti_join( auto const left_selected = matched.second.front(); auto const right_selected = matched.second.back(); - auto gather_map = - left_semi_anti_join(left_selected, right_selected, compare_nulls, stream); + auto gather_map = left_semi_anti_join(kind, left_selected, right_selected, compare_nulls, stream); auto const left_updated = scatter_columns(left_selected, left_on, left); return cudf::detail::gather(left_updated, @@ -224,8 +222,14 @@ std::unique_ptr left_semi_join(cudf::table_view const& left, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::left_semi_anti_join( - left, right, left_on, right_on, compare_nulls, rmm::cuda_stream_default, mr); + return detail::left_semi_anti_join(detail::join_kind::LEFT_SEMI_JOIN, + left, + right, + left_on, + right_on, + compare_nulls, + rmm::cuda_stream_default, + mr); } std::unique_ptr> left_semi_join( @@ -235,8 +239,8 @@ std::unique_ptr> left_semi_join( rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::left_semi_anti_join( - left, right, compare_nulls, rmm::cuda_stream_default, mr); + return detail::left_semi_anti_join( + detail::join_kind::LEFT_SEMI_JOIN, left, right, compare_nulls, rmm::cuda_stream_default, mr); } std::unique_ptr left_anti_join(cudf::table_view const& left, @@ -247,8 +251,14 @@ std::unique_ptr left_anti_join(cudf::table_view const& left, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::left_semi_anti_join( - left, right, left_on, right_on, compare_nulls, rmm::cuda_stream_default, mr); + return detail::left_semi_anti_join(detail::join_kind::LEFT_ANTI_JOIN, + left, + right, + left_on, + right_on, + compare_nulls, + rmm::cuda_stream_default, + mr); } std::unique_ptr> left_anti_join( @@ -258,8 +268,8 @@ std::unique_ptr> left_anti_join( rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::left_semi_anti_join( - left, right, compare_nulls, rmm::cuda_stream_default, mr); + return detail::left_semi_anti_join( + detail::join_kind::LEFT_ANTI_JOIN, left, right, compare_nulls, rmm::cuda_stream_default, mr); } } // namespace cudf From 26f18ef59f748e8674c3efc93386e27045ed606c Mon Sep 17 00:00:00 2001 From: davidwendt Date: Thu, 29 Jul 2021 19:12:54 -0400 Subject: [PATCH 2/4] make is_trivial_join utility not inline --- cpp/src/join/join.cu | 23 ++++++++++++++++++++++- cpp/src/join/join_common_utils.hpp | 25 ++----------------------- 2 files changed, 24 insertions(+), 24 deletions(-) diff --git a/cpp/src/join/join.cu b/cpp/src/join/join.cu index 526edbf6903..b43e3d1f6a9 100644 --- a/cpp/src/join/join.cu +++ b/cpp/src/join/join.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -26,6 +26,27 @@ namespace cudf { namespace detail { +// declared in join/join_common_utils.hpp +bool is_trivial_join(table_view const& left, table_view const& right, join_kind join_type) +{ + // If there is nothing to join, then send empty table with all columns + if (left.is_empty() || right.is_empty()) { return true; } + + // If left join and the left table is empty, return immediately + if ((join_kind::LEFT_JOIN == join_type) && (0 == left.num_rows())) { return true; } + + // If Inner Join and either table is empty, return immediately + if ((join_kind::INNER_JOIN == join_type) && ((0 == left.num_rows()) || (0 == right.num_rows()))) { + return true; + } + + // If left semi join (contains) and right table is empty, + // return immediately + if ((join_kind::LEFT_SEMI_JOIN == join_type) && (0 == right.num_rows())) { return true; } + + return false; +} + std::pair>, std::unique_ptr>> inner_join(table_view const& left_input, diff --git a/cpp/src/join/join_common_utils.hpp b/cpp/src/join/join_common_utils.hpp index d2337e28ed4..d2541b006a7 100644 --- a/cpp/src/join/join_common_utils.hpp +++ b/cpp/src/join/join_common_utils.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -19,8 +19,6 @@ #include #include -#include - #include #include @@ -49,26 +47,7 @@ using row_equality = cudf::row_equality_comparator; enum class join_kind { INNER_JOIN, LEFT_JOIN, FULL_JOIN, LEFT_SEMI_JOIN, LEFT_ANTI_JOIN }; -inline bool is_trivial_join(table_view const& left, table_view const& right, join_kind join_type) -{ - // If there is nothing to join, then send empty table with all columns - if (left.is_empty() || right.is_empty()) { return true; } - - // If left join and the left table is empty, return immediately - if ((join_kind::LEFT_JOIN == join_type) && (0 == left.num_rows())) { return true; } - - // If Inner Join and either table is empty, return immediately - if ((join_kind::INNER_JOIN == join_type) && ((0 == left.num_rows()) || (0 == right.num_rows()))) { - return true; - } - - // If left semi join (contains) and right table is empty, - // return immediately - if ((join_kind::LEFT_SEMI_JOIN == join_type) && (0 == right.num_rows())) { return true; } - - return false; -} +bool is_trivial_join(table_view const& left, table_view const& right, join_kind join_type); } // namespace detail - } // namespace cudf From dcc918ea94c9852371de5025d03a1722915c1aa7 Mon Sep 17 00:00:00 2001 From: davidwendt Date: Fri, 30 Jul 2021 10:38:18 -0400 Subject: [PATCH 3/4] use column instead of iterator for detail::gather call --- cpp/src/join/semi_join.cu | 19 +++++++++++++++---- 1 file changed, 15 insertions(+), 4 deletions(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 788b489d4ca..69a7b8c722b 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -19,7 +19,8 @@ #include #include -#include +#include +#include #include #include #include @@ -201,13 +202,23 @@ std::unique_ptr left_semi_anti_join( auto const left_selected = matched.second.front(); auto const right_selected = matched.second.back(); - auto gather_map = left_semi_anti_join(kind, left_selected, right_selected, compare_nulls, stream); + auto gather_vector = + left_semi_anti_join(kind, left_selected, right_selected, compare_nulls, stream); + + // wrapping the device vector with a column view allows calling the non-iterator + // version of detail::gather, improving compile time by 10% and reducing the + // object file size by 2.2x without affecting performance + auto gather_map = column_view(data_type{type_id::INT32}, + static_cast(gather_vector->size()), + gather_vector->data(), + nullptr, + 0); auto const left_updated = scatter_columns(left_selected, left_on, left); return cudf::detail::gather(left_updated, - gather_map->begin(), - gather_map->end(), + gather_map, out_of_bounds_policy::DONT_CHECK, + negative_index_policy::NOT_ALLOWED, stream, mr); } From 2270b9035b73a3ef33d631aba332d896f1ba5614 Mon Sep 17 00:00:00 2001 From: davidwendt Date: Fri, 30 Jul 2021 15:42:55 -0400 Subject: [PATCH 4/4] move more inline utilities to join_utils.cu --- cpp/CMakeLists.txt | 1 + cpp/src/join/join.cu | 21 ---- cpp/src/join/join_common_utils.cuh | 134 ++++--------------------- cpp/src/join/join_utils.cu | 155 +++++++++++++++++++++++++++++ 4 files changed, 178 insertions(+), 133 deletions(-) create mode 100644 cpp/src/join/join_utils.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 5c05a58b448..e1e5b7e624a 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -305,6 +305,7 @@ add_library(cudf src/join/cross_join.cu src/join/hash_join.cu src/join/join.cu + src/join/join_utils.cu src/join/semi_join.cu src/lists/contains.cu src/lists/combine/concatenate_list_elements.cu diff --git a/cpp/src/join/join.cu b/cpp/src/join/join.cu index b43e3d1f6a9..740431b8563 100644 --- a/cpp/src/join/join.cu +++ b/cpp/src/join/join.cu @@ -26,27 +26,6 @@ namespace cudf { namespace detail { -// declared in join/join_common_utils.hpp -bool is_trivial_join(table_view const& left, table_view const& right, join_kind join_type) -{ - // If there is nothing to join, then send empty table with all columns - if (left.is_empty() || right.is_empty()) { return true; } - - // If left join and the left table is empty, return immediately - if ((join_kind::LEFT_JOIN == join_type) && (0 == left.num_rows())) { return true; } - - // If Inner Join and either table is empty, return immediately - if ((join_kind::INNER_JOIN == join_type) && ((0 == left.num_rows()) || (0 == right.num_rows()))) { - return true; - } - - // If left semi join (contains) and right table is empty, - // return immediately - if ((join_kind::LEFT_SEMI_JOIN == join_type) && (0 == right.num_rows())) { return true; } - - return false; -} - std::pair>, std::unique_ptr>> inner_join(table_view const& left_input, diff --git a/cpp/src/join/join_common_utils.cuh b/cpp/src/join/join_common_utils.cuh index 2b1c870bea1..d5c23b1d612 100644 --- a/cpp/src/join/join_common_utils.cuh +++ b/cpp/src/join/join_common_utils.cuh @@ -21,9 +21,7 @@ #include #include -#include -#include #include namespace cudf { @@ -31,7 +29,9 @@ namespace detail { /** * @brief Computes the trivial left join operation for the case when the - * right table is empty. In this case all the valid indices of the left table + * right table is empty. + * + * In this case all the valid indices of the left table * are returned with their corresponding right indices being set to * JoinNoneValue, i.e. -1. * @@ -41,21 +41,12 @@ namespace detail { * * @return Join output indices vector pair */ -inline std::pair>, - std::unique_ptr>> +std::pair>, + std::unique_ptr>> get_trivial_left_join_indices( table_view const& left, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) -{ - auto left_indices = std::make_unique>(left.num_rows(), stream, mr); - thrust::sequence(rmm::exec_policy(stream), left_indices->begin(), left_indices->end(), 0); - auto right_indices = - std::make_unique>(left.num_rows(), stream, mr); - thrust::uninitialized_fill( - rmm::exec_policy(stream), right_indices->begin(), right_indices->end(), JoinNoneValue); - return std::make_pair(std::move(left_indices), std::move(right_indices)); -} + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); // Convenient alias for a pair of unique pointers to device uvectors. using VectorPair = std::pair>, @@ -83,47 +74,11 @@ using VectorPair = std::pair>, * * @return A pair of vectors containing the concatenated output. */ -inline VectorPair concatenate_vector_pairs(VectorPair& a, - VectorPair& b, - rmm::cuda_stream_view stream) -{ - CUDF_EXPECTS((a.first->size() == a.second->size()), - "Mismatch between sizes of vectors in vector pair"); - CUDF_EXPECTS((b.first->size() == b.second->size()), - "Mismatch between sizes of vectors in vector pair"); - if (a.first->is_empty()) { - return std::move(b); - } else if (b.first->is_empty()) { - return std::move(a); - } - auto original_size = a.first->size(); - a.first->resize(a.first->size() + b.first->size(), stream); - a.second->resize(a.second->size() + b.second->size(), stream); - thrust::copy( - rmm::exec_policy(stream), b.first->begin(), b.first->end(), a.first->begin() + original_size); - thrust::copy(rmm::exec_policy(stream), - b.second->begin(), - b.second->end(), - a.second->begin() + original_size); - return std::move(a); -} - -/** - * @brief Device functor to determine if an index is contained in a range. - */ -template -struct valid_range { - T start, stop; - __host__ __device__ valid_range(const T begin, const T end) : start(begin), stop(end) {} - - __host__ __device__ __forceinline__ bool operator()(const T index) - { - return ((index >= start) && (index < stop)); - } -}; +VectorPair concatenate_vector_pairs(VectorPair& a, VectorPair& b, rmm::cuda_stream_view stream); /** * @brief Creates a table containing the complement of left join indices. + * * This table has two columns. The first one is filled with JoinNoneValue(-1) * and the second one contains values from 0 to right_table_row_count - 1 * excluding those found in the right_indices column. @@ -136,72 +91,27 @@ struct valid_range { * * @return Pair of vectors containing the left join indices complement */ -inline std::pair>, - std::unique_ptr>> +std::pair>, + std::unique_ptr>> get_left_join_indices_complement(std::unique_ptr>& right_indices, size_type left_table_row_count, size_type right_table_row_count, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - // Get array of indices that do not appear in right_indices - - // Vector allocated for unmatched result - auto right_indices_complement = - std::make_unique>(right_table_row_count, stream); - - // If left table is empty in a full join call then all rows of the right table - // should be represented in the joined indices. This is an optimization since - // if left table is empty and full join is called all the elements in - // right_indices will be JoinNoneValue, i.e. -1. This if path should - // produce exactly the same result as the else path but will be faster. - if (left_table_row_count == 0) { - thrust::sequence(rmm::exec_policy(stream), - right_indices_complement->begin(), - right_indices_complement->end(), - 0); - } else { - // Assume all the indices in invalid_index_map are invalid - auto invalid_index_map = - std::make_unique>(right_table_row_count, stream); - thrust::uninitialized_fill( - rmm::exec_policy(stream), invalid_index_map->begin(), invalid_index_map->end(), int32_t{1}); - - // Functor to check for index validity since left joins can create invalid indices - valid_range valid(0, right_table_row_count); + rmm::mr::device_memory_resource* mr); - // invalid_index_map[index_ptr[i]] = 0 for i = 0 to right_table_row_count - // Thus specifying that those locations are valid - thrust::scatter_if(rmm::exec_policy(stream), - thrust::make_constant_iterator(0), - thrust::make_constant_iterator(0) + right_indices->size(), - right_indices->begin(), // Index locations - right_indices->begin(), // Stencil - Check if index location is valid - invalid_index_map->begin(), // Output indices - valid); // Stencil Predicate - size_type begin_counter = static_cast(0); - size_type end_counter = static_cast(right_table_row_count); +/** + * @brief Device functor to determine if an index is contained in a range. + */ +template +struct valid_range { + T start, stop; + __host__ __device__ valid_range(const T begin, const T end) : start(begin), stop(end) {} - // Create list of indices that have been marked as invalid - size_type indices_count = thrust::copy_if(rmm::exec_policy(stream), - thrust::make_counting_iterator(begin_counter), - thrust::make_counting_iterator(end_counter), - invalid_index_map->begin(), - right_indices_complement->begin(), - thrust::identity()) - - right_indices_complement->begin(); - right_indices_complement->resize(indices_count, stream); + __host__ __device__ __forceinline__ bool operator()(const T index) + { + return ((index >= start) && (index < stop)); } - - auto left_invalid_indices = - std::make_unique>(right_indices_complement->size(), stream); - thrust::uninitialized_fill(rmm::exec_policy(stream), - left_invalid_indices->begin(), - left_invalid_indices->end(), - JoinNoneValue); - - return std::make_pair(std::move(left_invalid_indices), std::move(right_indices_complement)); -} +}; /** * @brief Adds a pair of indices to the shared memory cache diff --git a/cpp/src/join/join_utils.cu b/cpp/src/join/join_utils.cu new file mode 100644 index 00000000000..4aca4b4a9cf --- /dev/null +++ b/cpp/src/join/join_utils.cu @@ -0,0 +1,155 @@ +/* + * 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 + +namespace cudf { +namespace detail { + +bool is_trivial_join(table_view const& left, table_view const& right, join_kind join_type) +{ + // If there is nothing to join, then send empty table with all columns + if (left.is_empty() || right.is_empty()) { return true; } + + // If left join and the left table is empty, return immediately + if ((join_kind::LEFT_JOIN == join_type) && (0 == left.num_rows())) { return true; } + + // If Inner Join and either table is empty, return immediately + if ((join_kind::INNER_JOIN == join_type) && ((0 == left.num_rows()) || (0 == right.num_rows()))) { + return true; + } + + // If left semi join (contains) and right table is empty, + // return immediately + if ((join_kind::LEFT_SEMI_JOIN == join_type) && (0 == right.num_rows())) { return true; } + + return false; +} + +std::pair>, + std::unique_ptr>> +get_trivial_left_join_indices(table_view const& left, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + auto left_indices = std::make_unique>(left.num_rows(), stream, mr); + thrust::sequence(rmm::exec_policy(stream), left_indices->begin(), left_indices->end(), 0); + auto right_indices = + std::make_unique>(left.num_rows(), stream, mr); + thrust::uninitialized_fill( + rmm::exec_policy(stream), right_indices->begin(), right_indices->end(), JoinNoneValue); + return std::make_pair(std::move(left_indices), std::move(right_indices)); +} + +VectorPair concatenate_vector_pairs(VectorPair& a, VectorPair& b, rmm::cuda_stream_view stream) +{ + CUDF_EXPECTS((a.first->size() == a.second->size()), + "Mismatch between sizes of vectors in vector pair"); + CUDF_EXPECTS((b.first->size() == b.second->size()), + "Mismatch between sizes of vectors in vector pair"); + if (a.first->is_empty()) { + return std::move(b); + } else if (b.first->is_empty()) { + return std::move(a); + } + auto original_size = a.first->size(); + a.first->resize(a.first->size() + b.first->size(), stream); + a.second->resize(a.second->size() + b.second->size(), stream); + thrust::copy( + rmm::exec_policy(stream), b.first->begin(), b.first->end(), a.first->begin() + original_size); + thrust::copy(rmm::exec_policy(stream), + b.second->begin(), + b.second->end(), + a.second->begin() + original_size); + return std::move(a); +} + +std::pair>, + std::unique_ptr>> +get_left_join_indices_complement(std::unique_ptr>& right_indices, + size_type left_table_row_count, + size_type right_table_row_count, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + // Get array of indices that do not appear in right_indices + + // Vector allocated for unmatched result + auto right_indices_complement = + std::make_unique>(right_table_row_count, stream); + + // If left table is empty in a full join call then all rows of the right table + // should be represented in the joined indices. This is an optimization since + // if left table is empty and full join is called all the elements in + // right_indices will be JoinNoneValue, i.e. -1. This if path should + // produce exactly the same result as the else path but will be faster. + if (left_table_row_count == 0) { + thrust::sequence(rmm::exec_policy(stream), + right_indices_complement->begin(), + right_indices_complement->end(), + 0); + } else { + // Assume all the indices in invalid_index_map are invalid + auto invalid_index_map = + std::make_unique>(right_table_row_count, stream); + thrust::uninitialized_fill( + rmm::exec_policy(stream), invalid_index_map->begin(), invalid_index_map->end(), int32_t{1}); + + // Functor to check for index validity since left joins can create invalid indices + valid_range valid(0, right_table_row_count); + + // invalid_index_map[index_ptr[i]] = 0 for i = 0 to right_table_row_count + // Thus specifying that those locations are valid + thrust::scatter_if(rmm::exec_policy(stream), + thrust::make_constant_iterator(0), + thrust::make_constant_iterator(0) + right_indices->size(), + right_indices->begin(), // Index locations + right_indices->begin(), // Stencil - Check if index location is valid + invalid_index_map->begin(), // Output indices + valid); // Stencil Predicate + size_type begin_counter = static_cast(0); + size_type end_counter = static_cast(right_table_row_count); + + // Create list of indices that have been marked as invalid + size_type indices_count = thrust::copy_if(rmm::exec_policy(stream), + thrust::make_counting_iterator(begin_counter), + thrust::make_counting_iterator(end_counter), + invalid_index_map->begin(), + right_indices_complement->begin(), + thrust::identity()) - + right_indices_complement->begin(); + right_indices_complement->resize(indices_count, stream); + } + + auto left_invalid_indices = + std::make_unique>(right_indices_complement->size(), stream); + thrust::uninitialized_fill(rmm::exec_policy(stream), + left_invalid_indices->begin(), + left_invalid_indices->end(), + JoinNoneValue); + + return std::make_pair(std::move(left_invalid_indices), std::move(right_indices_complement)); +} + +} // namespace detail +} // namespace cudf