diff --git a/cpp/include/cudf/detail/null_mask.cuh b/cpp/include/cudf/detail/null_mask.cuh index 02c93e7e31e..753525128bb 100644 --- a/cpp/include/cudf/detail/null_mask.cuh +++ b/cpp/include/cudf/detail/null_mask.cuh @@ -98,16 +98,17 @@ __global__ void offset_bitmask_binop(Binop op, * @param stream CUDA stream used for device memory operations and kernel launches */ template -bitmask bitmask_binop(Binop op, - host_span masks, - host_span masks_begin_bits, - size_type mask_size_bits, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +std::pair bitmask_binop( + Binop op, + host_span masks, + host_span masks_begin_bits, + size_type mask_size_bits, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { auto dest_mask = rmm::device_buffer{bitmask_allocation_size_bytes(mask_size_bits), stream, mr}; - - auto valid_count = + auto null_count = + mask_size_bits - inplace_bitmask_binop(op, device_span(static_cast(dest_mask.data()), num_bitmask_words(mask_size_bits)), @@ -116,9 +117,8 @@ bitmask bitmask_binop(Binop op, mask_size_bits, stream, mr); - auto null_count = mask_size_bits - valid_count; - return bitmask{std::move(dest_mask), valid_count, null_count}; + return std::make_pair(std::move(dest_mask), null_count); } /** diff --git a/cpp/include/cudf/detail/null_mask.hpp b/cpp/include/cudf/detail/null_mask.hpp index 334d0339ca2..d2819e665df 100644 --- a/cpp/include/cudf/detail/null_mask.hpp +++ b/cpp/include/cudf/detail/null_mask.hpp @@ -23,9 +23,6 @@ #include namespace cudf { - -struct bitmask; - namespace detail { /** @@ -117,29 +114,32 @@ rmm::device_buffer copy_bitmask( * * @param stream CUDA stream used for device memory operations and kernel launches */ -bitmask bitmask_and(host_span masks, - host_span masks_begin_bits, - size_type mask_size_bits, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::pair bitmask_and( + host_span masks, + host_span masks_begin_bits, + size_type mask_size_bits, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** * @copydoc cudf::bitmask_and * * @param[in] stream CUDA stream used for device memory operations and kernel launches. */ -bitmask bitmask_and(table_view const& view, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::pair bitmask_and( + table_view const& view, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** * @copydoc cudf::bitmask_or * * @param[in] stream CUDA stream used for device memory operations and kernel launches. */ -bitmask bitmask_or(table_view const& view, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::pair bitmask_or( + table_view const& view, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** * @brief Performs a bitwise AND of the specified bitmasks, diff --git a/cpp/include/cudf/null_mask.hpp b/cpp/include/cudf/null_mask.hpp index 1ea3160161a..4b887b20049 100644 --- a/cpp/include/cudf/null_mask.hpp +++ b/cpp/include/cudf/null_mask.hpp @@ -24,15 +24,6 @@ namespace cudf { -/** - * @brief Bitmask output type. - */ -struct bitmask { - rmm::device_buffer mask; ///< Resulting bitmask - size_type num_set_bits; ///< Number of set bits - size_type num_unset_bits; ///< Number of unset bits -}; - /** * @addtogroup column_nullmask * @{ @@ -211,32 +202,32 @@ rmm::device_buffer copy_bitmask( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Returns a struct of bitwise AND of the bitmasks of columns of a table, - * count of valid bits and count of null bits + * @brief Returns a pair of bitwise AND of the bitmasks of columns of a table and count of null bits * * If any of the columns isn't nullable, it is considered all valid. * If no column in the table is nullable, an empty bitmask is returned. * * @param view The table of columns * @param mr Device memory resource used to allocate the returned device_buffer - * @return A struct of resulting bitmask, count of valid bits and count of null bits + * @return A pair of resulting bitmask and count of null bits */ -bitmask bitmask_and(table_view const& view, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::pair bitmask_and( + table_view const& view, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Returns a struct of bitwise OR of the bitmasks of columns of a table, - * count of valid bits and count of null bits + * @brief Returns a pair of bitwise OR of the bitmasks of columns of a table and count of null bits * * If any of the columns isn't nullable, it is considered all valid. * If no column in the table is nullable, an empty bitmask is returned. * * @param view The table of columns * @param mr Device memory resource used to allocate the returned device_buffer - * @return rmm::device_buffer Output bitmask + * @return A pair of resulting bitmask and count of null bits */ -bitmask bitmask_or(table_view const& view, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::pair bitmask_or( + table_view const& view, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @} */ // end of group } // namespace cudf diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 794d0676f8a..b9ed95daf1b 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -392,13 +392,9 @@ std::unique_ptr make_fixed_width_column_for_output(column_view const& lh if (binops::is_null_dependent(op)) { return make_fixed_width_column(output_type, rhs.size(), mask_state::ALL_VALID, stream, mr); } else { - auto bitmask_output = cudf::detail::bitmask_and(table_view({lhs, rhs}), stream, mr); - return make_fixed_width_column(output_type, - lhs.size(), - std::move(bitmask_output.mask), - bitmask_output.num_unset_bits, - stream, - mr); + auto [new_mask, null_count] = cudf::detail::bitmask_and(table_view({lhs, rhs}), stream, mr); + return make_fixed_width_column( + output_type, lhs.size(), std::move(new_mask), null_count, stream, mr); } }; @@ -803,13 +799,9 @@ std::unique_ptr binary_operation(column_view const& lhs, CUDF_EXPECTS((lhs.size() == rhs.size()), "Column sizes don't match"); - auto bitmask_output = bitmask_and(table_view({lhs, rhs}), stream, mr); - auto out = make_fixed_width_column(output_type, - lhs.size(), - std::move(bitmask_output.mask), - bitmask_output.num_unset_bits, - stream, - mr); + auto [new_mask, null_count] = bitmask_and(table_view({lhs, rhs}), stream, mr); + auto out = + make_fixed_width_column(output_type, lhs.size(), std::move(new_mask), null_count, stream, mr); // Check for 0 sized data if (lhs.is_empty() or rhs.is_empty()) return out; diff --git a/cpp/src/bitmask/null_mask.cu b/cpp/src/bitmask/null_mask.cu index 09560396c21..1cd3def61ac 100644 --- a/cpp/src/bitmask/null_mask.cu +++ b/cpp/src/bitmask/null_mask.cu @@ -356,11 +356,11 @@ void inplace_bitmask_and(device_span dest_mask, } // Bitwise AND of the masks -bitmask bitmask_and(host_span masks, - host_span begin_bits, - size_type mask_size, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::pair bitmask_and(host_span masks, + host_span begin_bits, + size_type mask_size, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { return bitmask_binop( [] __device__(bitmask_type left, bitmask_type right) { return left & right; }, @@ -372,14 +372,14 @@ bitmask bitmask_and(host_span masks, } // Returns the bitwise AND of the null masks of all columns in the table view -bitmask bitmask_and(table_view const& view, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::pair bitmask_and(table_view const& view, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); rmm::device_buffer null_mask{0, stream, mr}; if (view.num_rows() == 0 or view.num_columns() == 0) { - return bitmask{std::move(null_mask), 0, 0}; + return std::make_pair(std::move(null_mask), 0); } std::vector masks; @@ -401,18 +401,18 @@ bitmask bitmask_and(table_view const& view, mr); } - return bitmask{std::move(null_mask), 0, 0}; + return std::make_pair(std::move(null_mask), 0); } // Returns the bitwise OR of the null masks of all columns in the table view -bitmask bitmask_or(table_view const& view, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::pair bitmask_or(table_view const& view, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); rmm::device_buffer null_mask{0, stream, mr}; if (view.num_rows() == 0 or view.num_columns() == 0) { - return bitmask{std::move(null_mask), 0, 0}; + return std::make_pair(std::move(null_mask), 0); } std::vector masks; @@ -434,7 +434,7 @@ bitmask bitmask_or(table_view const& view, mr); } - return bitmask{std::move(null_mask), 0, 0}; + return std::make_pair(std::move(null_mask), 0); } /** @@ -506,12 +506,14 @@ rmm::device_buffer copy_bitmask(column_view const& view, rmm::mr::device_memory_ return detail::copy_bitmask(view, rmm::cuda_stream_default, mr); } -bitmask bitmask_and(table_view const& view, rmm::mr::device_memory_resource* mr) +std::pair bitmask_and(table_view const& view, + rmm::mr::device_memory_resource* mr) { return detail::bitmask_and(view, rmm::cuda_stream_default, mr); } -bitmask bitmask_or(table_view const& view, rmm::mr::device_memory_resource* mr) +std::pair bitmask_or(table_view const& view, + rmm::mr::device_memory_resource* mr) { return detail::bitmask_or(view, rmm::cuda_stream_default, mr); } diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index fef720da174..7d66daf226e 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -364,9 +364,9 @@ std::unique_ptr add_calendrical_months(column_view const& timestamp_colu stream, mr); - auto bitmask_output = + auto [output_null_mask, null_count] = cudf::detail::bitmask_and(table_view{{timestamp_column, months_column}}, stream, mr); - output->set_null_mask(std::move(bitmask_output.mask), bitmask_output.num_unset_bits); + output->set_null_mask(std::move(output_null_mask), null_count); return output; } diff --git a/cpp/src/groupby/hash/groupby.cu b/cpp/src/groupby/hash/groupby.cu index c29fc272078..1b9b2a196c3 100644 --- a/cpp/src/groupby/hash/groupby.cu +++ b/cpp/src/groupby/hash/groupby.cu @@ -390,7 +390,7 @@ void sparse_to_dense_results(table_view const& keys, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto row_bitmask = bitmask_and(keys, stream, rmm::mr::get_current_device_resource()).mask; + auto row_bitmask = bitmask_and(keys, stream, rmm::mr::get_current_device_resource()).first; bool skip_key_rows_with_nulls = keys_have_nulls and include_null_keys == null_policy::EXCLUDE; bitmask_type const* row_bitmask_ptr = skip_key_rows_with_nulls ? static_cast(row_bitmask.data()) : nullptr; @@ -502,7 +502,7 @@ void compute_single_pass_aggs(table_view const& keys, bool skip_key_rows_with_nulls = keys_have_nulls and include_null_keys == null_policy::EXCLUDE; auto row_bitmask = - skip_key_rows_with_nulls ? cudf::detail::bitmask_and(keys, stream).mask : rmm::device_buffer{}; + skip_key_rows_with_nulls ? cudf::detail::bitmask_and(keys, stream).first : rmm::device_buffer{}; thrust::for_each_n( rmm::exec_policy(stream), thrust::make_counting_iterator(0), diff --git a/cpp/src/groupby/sort/aggregate.cpp b/cpp/src/groupby/sort/aggregate.cpp index b263cbb9732..234bb447761 100644 --- a/cpp/src/groupby/sort/aggregate.cpp +++ b/cpp/src/groupby/sort/aggregate.cpp @@ -538,10 +538,7 @@ void aggregate_result_functor::operator()(aggregation con */ auto column_view_with_common_nulls(column_view const& column_0, column_view const& column_1) { - auto bitmask_output = cudf::bitmask_and(table_view{{column_0, column_1}}); - auto new_nullmask = std::move(bitmask_output.mask); - auto null_count = bitmask_output.num_unset_bits; - + auto [new_nullmask, null_count] = cudf::bitmask_and(table_view{{column_0, column_1}}); if (null_count == 0) { return std::make_tuple(std::move(new_nullmask), column_0, column_1); } auto column_view_with_new_nullmask = [](auto const& col, void* nullmask, auto null_count) { return column_view(col.type(), diff --git a/cpp/src/groupby/sort/sort_helper.cu b/cpp/src/groupby/sort/sort_helper.cu index be26d9e52af..7adb4ccec76 100644 --- a/cpp/src/groupby/sort/sort_helper.cu +++ b/cpp/src/groupby/sort/sort_helper.cu @@ -276,13 +276,10 @@ column_view sort_groupby_helper::keys_bitmask_column(rmm::cuda_stream_view strea { if (_keys_bitmask_column) return _keys_bitmask_column->view(); - auto bitmask_output = cudf::detail::bitmask_and(_keys, stream); + auto [row_bitmask, null_count] = cudf::detail::bitmask_and(_keys, stream); - _keys_bitmask_column = make_numeric_column(data_type(type_id::INT8), - _keys.num_rows(), - std::move(bitmask_output.mask), - bitmask_output.num_unset_bits, - stream); + _keys_bitmask_column = make_numeric_column( + data_type(type_id::INT8), _keys.num_rows(), std::move(row_bitmask), null_count, stream); auto keys_bitmask_view = _keys_bitmask_column->mutable_view(); using T = id_to_type; diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index affde8228ce..d5065278afc 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -92,7 +92,7 @@ void build_join_hash_table(cudf::table_view const& build, hash_table.insert(iter, iter + build_table_num_rows, stream.value()); } else { thrust::counting_iterator stencil(0); - auto const row_bitmask = cudf::detail::bitmask_and(build, stream).mask; + auto const row_bitmask = cudf::detail::bitmask_and(build, stream).first; row_is_valid pred{static_cast(row_bitmask.data())}; // insert valid rows diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index b6433578077..5b5dd418a97 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -97,7 +97,7 @@ std::unique_ptr> left_semi_anti_join( // 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{} - : cudf::detail::bitmask_and(right_flattened_keys, stream).mask; + : cudf::detail::bitmask_and(right_flattened_keys, stream).first; // skip rows that are null here. thrust::for_each_n( rmm::exec_policy(stream), diff --git a/cpp/src/strings/repeat_strings.cu b/cpp/src/strings/repeat_strings.cu index beaa68b9cc5..458f3ed885c 100644 --- a/cpp/src/strings/repeat_strings.cu +++ b/cpp/src/strings/repeat_strings.cu @@ -319,14 +319,14 @@ std::unique_ptr repeat_strings(strings_column_view const& input, // We generate new bitmask by AND of the input columns' bitmasks. // Note that if the input columns are nullable, the output column will also be nullable (which may // not have nulls). - auto bitmask_out = + auto [null_mask, null_count] = cudf::detail::bitmask_and(table_view{{input.parent(), repeat_times}}, stream, mr); return make_strings_column(strings_count, std::move(offsets_column), std::move(chars_column), - bitmask_out.num_unset_bits, - std::move(bitmask_out.mask)); + null_count, + std::move(null_mask)); } std::pair, int64_t> repeat_strings_output_sizes( diff --git a/cpp/src/structs/utilities.cpp b/cpp/src/structs/utilities.cpp index 7628ea51f15..d4e2f48feba 100644 --- a/cpp/src/structs/utilities.cpp +++ b/cpp/src/structs/utilities.cpp @@ -379,15 +379,14 @@ std::tuple> superimpose_paren // and the _null_mask(). It would be better to AND the bits from the beginning, and apply // offset() uniformly. // Alternatively, one could construct a big enough buffer, and use inplace_bitwise_and. - auto bitmask_output = cudf::detail::bitmask_and(parent_child_null_masks, - std::vector{0, 0}, - child.offset() + child.size(), - stream, - mr); - ret_validity_buffers.push_back(std::move(bitmask_output.mask)); + auto [new_mask, null_count] = cudf::detail::bitmask_and(parent_child_null_masks, + std::vector{0, 0}, + child.offset() + child.size(), + stream, + mr); + ret_validity_buffers.push_back(std::move(new_mask)); return std::make_pair( - reinterpret_cast(ret_validity_buffers.back().data()), - bitmask_output.num_unset_bits); + reinterpret_cast(ret_validity_buffers.back().data()), null_count); }(); return cudf::column_view( diff --git a/cpp/tests/bitmask/bitmask_tests.cpp b/cpp/tests/bitmask/bitmask_tests.cpp index 72853f1f15a..c7ae6e12366 100644 --- a/cpp/tests/bitmask/bitmask_tests.cpp +++ b/cpp/tests/bitmask/bitmask_tests.cpp @@ -545,29 +545,25 @@ TEST_F(MergeBitmaskTest, TestBitmaskAnd) auto const input2 = cudf::table_view({bools_col1, bools_col2}); auto const input3 = cudf::table_view({bools_col1, bools_col2, bools_col3}); - auto result1 = cudf::bitmask_and(input1); - auto result2 = cudf::bitmask_and(input2); - auto result3 = cudf::bitmask_and(input3); + auto [result1_mask, result1_null_count] = cudf::bitmask_and(input1); + auto [result2_mask, result2_null_count] = cudf::bitmask_and(input2); + auto [result3_mask, result3_null_count] = cudf::bitmask_and(input3); - constexpr cudf::size_type gold_valid_count = 2; - constexpr cudf::size_type gold_null_count = 3; + constexpr cudf::size_type gold_null_count = 3; - EXPECT_EQ(result1.num_set_bits, 0); - EXPECT_EQ(result1.num_unset_bits, 0); - EXPECT_EQ(result2.num_set_bits, gold_valid_count); - EXPECT_EQ(result2.num_unset_bits, gold_null_count); - EXPECT_EQ(result3.num_set_bits, gold_valid_count); - EXPECT_EQ(result3.num_unset_bits, gold_null_count); + EXPECT_EQ(result1_null_count, 0); + EXPECT_EQ(result2_null_count, gold_null_count); + EXPECT_EQ(result3_null_count, gold_null_count); auto odd_indices = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i % 2; }); auto odd = cudf::test::detail::make_null_mask(odd_indices, odd_indices + input2.num_rows()); - EXPECT_EQ(nullptr, result1.mask.data()); + EXPECT_EQ(nullptr, result1_mask.data()); CUDF_TEST_EXPECT_EQUAL_BUFFERS( - result2.mask.data(), odd.data(), cudf::num_bitmask_words(input2.num_rows())); + result2_mask.data(), odd.data(), cudf::num_bitmask_words(input2.num_rows())); CUDF_TEST_EXPECT_EQUAL_BUFFERS( - result3.mask.data(), odd.data(), cudf::num_bitmask_words(input2.num_rows())); + result3_mask.data(), odd.data(), cudf::num_bitmask_words(input2.num_rows())); } TEST_F(MergeBitmaskTest, TestBitmaskOr) @@ -580,29 +576,23 @@ TEST_F(MergeBitmaskTest, TestBitmaskOr) auto const input2 = cudf::table_view({bools_col1, bools_col2}); auto const input3 = cudf::table_view({bools_col1, bools_col2, bools_col3}); - auto result1 = cudf::bitmask_or(input1); - auto result2 = cudf::bitmask_or(input2); - auto result3 = cudf::bitmask_or(input3); + auto [result1_mask, result1_null_count] = cudf::bitmask_or(input1); + auto [result2_mask, result2_null_count] = cudf::bitmask_or(input2); + auto [result3_mask, result3_null_count] = cudf::bitmask_or(input3); - constexpr cudf::size_type gold_valid_count = 4; - constexpr cudf::size_type gold_null_count = 1; - - EXPECT_EQ(result1.num_set_bits, 0); - EXPECT_EQ(result1.num_unset_bits, 0); - EXPECT_EQ(result2.num_set_bits, gold_valid_count); - EXPECT_EQ(result2.num_unset_bits, gold_null_count); - EXPECT_EQ(result3.num_set_bits, 0); - EXPECT_EQ(result3.num_unset_bits, 0); + EXPECT_EQ(result1_null_count, 0); + EXPECT_EQ(result2_null_count, 1); + EXPECT_EQ(result3_null_count, 0); auto all_but_index3 = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 3; }); auto null3 = cudf::test::detail::make_null_mask(all_but_index3, all_but_index3 + input2.num_rows()); - EXPECT_EQ(nullptr, result1.mask.data()); + EXPECT_EQ(nullptr, result1_mask.data()); CUDF_TEST_EXPECT_EQUAL_BUFFERS( - result2.mask.data(), null3.data(), cudf::num_bitmask_words(input2.num_rows())); - EXPECT_EQ(nullptr, result3.mask.data()); + result2_mask.data(), null3.data(), cudf::num_bitmask_words(input2.num_rows())); + EXPECT_EQ(nullptr, result3_mask.data()); } CUDF_TEST_PROGRAM_MAIN() diff --git a/java/src/main/native/src/ColumnViewJni.cpp b/java/src/main/native/src/ColumnViewJni.cpp index 679f66887a7..5ae9fd03063 100644 --- a/java/src/main/native/src/ColumnViewJni.cpp +++ b/java/src/main/native/src/ColumnViewJni.cpp @@ -1493,13 +1493,13 @@ JNIEXPORT jlong JNICALL Java_ai_rapids_cudf_ColumnView_bitwiseMergeAndSetValidit cudf::binary_operator op = static_cast(bin_op); switch (op) { case cudf::binary_operator::BITWISE_AND: { - auto bitmask_output = cudf::bitmask_and(*input_table); - copy->set_null_mask(std::move(bitmask_output.mask), bitmask_output.num_unset_bits); + auto [new_bitmask, null_count] = cudf::bitmask_and(*input_table); + copy->set_null_mask(std::move(new_bitmask), null_count); break; } case cudf::binary_operator::BITWISE_OR: { - auto bitmask_output = cudf::bitmask_or(*input_table); - copy->set_null_mask(std::move(bitmask_output.mask), bitmask_output.num_unset_bits); + auto [new_bitmask, null_count] = cudf::bitmask_or(*input_table); + copy->set_null_mask(std::move(new_bitmask), null_count); break; } default: JNI_THROW_NEW(env, cudf::jni::ILLEGAL_ARG_CLASS, "Unsupported merge operation", 0);