diff --git a/cpp/benchmarks/bitmask/bitmask_and.cpp b/cpp/benchmarks/bitmask/bitmask_and.cpp index 00c11e200c05..4d48e108de39 100644 --- a/cpp/benchmarks/bitmask/bitmask_and.cpp +++ b/cpp/benchmarks/bitmask/bitmask_and.cpp @@ -44,12 +44,12 @@ auto setup_masks(nvbench::state& state) std::exclusive_scan(segments.begin(), segments.end(), segments.begin(), 0); // Create masks - std::vector> masks; + std::vector> masks; std::vector mask_pointers; masks.reserve(num_masks); std::generate_n(std::back_inserter(masks), num_masks, [mask_size_bits, seed, &mask_pointers]() { auto mask_pair = create_random_null_mask(mask_size_bits, null_probability, seed); - mask_pointers.push_back(reinterpret_cast(mask_pair.first.data())); + mask_pointers.push_back(mask_pair.first.data()); return std::move(mask_pair.first); }); diff --git a/cpp/benchmarks/bitmask/mask_to_bools.cpp b/cpp/benchmarks/bitmask/mask_to_bools.cpp index 7553bc68779d..30a782dd3933 100644 --- a/cpp/benchmarks/bitmask/mask_to_bools.cpp +++ b/cpp/benchmarks/bitmask/mask_to_bools.cpp @@ -13,7 +13,7 @@ static void bench_mask_to_bools(nvbench::state& state) { auto const num_rows = static_cast(state.get_int64("num_rows")); auto const mask = cudf::create_null_mask(num_rows, cudf::mask_state::ALL_VALID); - auto const bitmask = reinterpret_cast(mask.data()); + auto const bitmask = mask.data(); state.set_cuda_stream(nvbench::make_cuda_stream_view(cudf::get_default_stream().get())); state.add_global_memory_reads(cudf::bitmask_allocation_size_bytes(num_rows)); diff --git a/cpp/benchmarks/bitmask/set_null_mask.cpp b/cpp/benchmarks/bitmask/set_null_mask.cpp index 1b0114bbb02e..5d1f922476be 100644 --- a/cpp/benchmarks/bitmask/set_null_mask.cpp +++ b/cpp/benchmarks/bitmask/set_null_mask.cpp @@ -39,14 +39,14 @@ auto generate_test_data(cudf::size_type num_masks, auto valids = thrust::host_vector(num_masks, true); - std::vector> masks; + std::vector> masks; masks.reserve(num_masks); std::vector masks_ptr; masks_ptr.reserve(num_masks); for (cudf::size_type i = 0; i < num_masks; ++i) { auto& last = masks.emplace_back(cudf::create_null_mask(mask_size, cudf::mask_state::UNINITIALIZED)); - masks_ptr.emplace_back(reinterpret_cast(last.data())); + masks_ptr.emplace_back(last.data()); } return std::make_tuple(std::move(begin_bits), @@ -61,19 +61,19 @@ auto generate_test_data(cudf::size_type num_masks, void BM_setnullmask(nvbench::state& state) { auto const mask_size = static_cast(state.get_int64("mask_size")); - cuda::device_buffer mask = + cuda::device_buffer mask = cudf::create_null_mask(mask_size, cudf::mask_state::UNINITIALIZED); auto begin = 0, end = mask_size; state.set_cuda_stream(nvbench::make_cuda_stream_view(cudf::get_default_stream().get())); auto const mem_stats_logger = cudf::memory_stats_logger(); - state.exec( - nvbench::exec_tag::sync | nvbench::exec_tag::timer, [&](nvbench::launch& launch, auto& timer) { - timer.start(); - cudf::set_null_mask(reinterpret_cast(mask.data()), begin, end, true); - timer.stop(); - }); + state.exec(nvbench::exec_tag::sync | nvbench::exec_tag::timer, + [&](nvbench::launch& launch, auto& timer) { + timer.start(); + cudf::set_null_mask(mask.data(), begin, end, true); + timer.stop(); + }); state.add_buffer_size( mem_stats_logger.peak_memory_usage(), "peak_memory_usage", "peak_memory_usage"); diff --git a/cpp/benchmarks/common/generate_input.cu b/cpp/benchmarks/common/generate_input.cu index 01f8cefaec0f..8cd179161003 100644 --- a/cpp/benchmarks/common/generate_input.cu +++ b/cpp/benchmarks/common/generate_input.cu @@ -527,7 +527,7 @@ std::unique_ptr create_random_utf8_string_column(data_profile cons auto [result_bitmask, null_count] = profile.get_null_probability().has_value() ? cudf::bools_to_mask(cudf::device_span(null_mask), stream) - : std::pair{std::make_unique>( + : std::pair{std::make_unique>( cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)), 0}; @@ -640,7 +640,7 @@ std::unique_ptr create_random_column(data_profile const& profile, auto [result_bitmask, null_count] = profile.get_null_probability().has_value() ? cudf::bools_to_mask(cudf::device_span(null_mask)) - : std::pair{std::make_unique>( + : std::pair{std::make_unique>( cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)), 0}; @@ -720,7 +720,7 @@ std::unique_ptr create_random_column(data_profi return cudf::bools_to_mask(cudf::device_span(valids), cudf::get_default_stream()); } - return std::pair{std::make_unique>( + return std::pair{std::make_unique>( cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)), 0}; }(); @@ -817,7 +817,7 @@ std::unique_ptr create_random_column(data_profile auto [null_mask, null_count] = profile.get_null_probability().has_value() ? cudf::bools_to_mask(cudf::device_span(valids)) - : std::pair{std::make_unique>( + : std::pair{std::make_unique>( cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)), 0}; @@ -1099,7 +1099,7 @@ std::unique_ptr create_string_column(cudf::size_type num_rows, return std::move(table->release().front()); } -std::pair, cudf::size_type> create_random_null_mask( +std::pair, cudf::size_type> create_random_null_mask( cudf::size_type size, std::optional null_probability, unsigned seed) { if (not null_probability.has_value()) { diff --git a/cpp/benchmarks/common/generate_input.hpp b/cpp/benchmarks/common/generate_input.hpp index 527d445d77ef..67ba99f49878 100644 --- a/cpp/benchmarks/common/generate_input.hpp +++ b/cpp/benchmarks/common/generate_input.hpp @@ -674,5 +674,5 @@ std::vector mix_dtypes(std::pair co * @param seed Optional, seed for the pseudo-random engine * @return null mask device buffer with random null mask data and null count */ -std::pair, cudf::size_type> create_random_null_mask( +std::pair, cudf::size_type> create_random_null_mask( cudf::size_type size, std::optional null_probability = std::nullopt, unsigned seed = 1); diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index 1505dd21c2e3..020e9c1b94b7 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -407,7 +407,7 @@ bool is_supported_operation(data_type out, data_type lhs, data_type rhs, binary_ * @param mr Device memory resource used to allocate the returned valid mask * @return Computed validity mask */ -std::pair, size_type> scalar_col_valid_mask_and( +std::pair, size_type> scalar_col_valid_mask_and( column_view const& col, scalar const& s, cuda::stream_ref stream = cudf::get_default_stream(), diff --git a/cpp/include/cudf/column/column.hpp b/cpp/include/cudf/column/column.hpp index 8881b3f01427..093fc4ecfb3a 100644 --- a/cpp/include/cudf/column/column.hpp +++ b/cpp/include/cudf/column/column.hpp @@ -74,7 +74,7 @@ class column { */ template () or cudf::is_chrono())> column(rmm::device_uvector&& other, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, size_type null_count) : _type{cudf::data_type{cudf::type_to_id()}}, _size{[&]() { @@ -105,7 +105,7 @@ class column { * @param null_count Optional, the count of null elements. * @param children Optional, vector of child columns */ - template > + template > column(data_type dtype, size_type size, B1&& data, @@ -168,7 +168,8 @@ class column { * `new_null_count` is 0. * @param new_null_count The count of null elements. */ - void set_null_mask(cuda::device_buffer&& new_null_mask, size_type new_null_count); + void set_null_mask(cuda::device_buffer&& new_null_mask, + size_type new_null_count); /** * @brief Sets the column's null value indicator bitmask to `new_null_mask`. @@ -182,7 +183,7 @@ class column { * @param stream The stream on which to perform the allocation and copy. Uses the default CUDF * stream if none is specified. */ - void set_null_mask(cuda::device_buffer const& new_null_mask, + void set_null_mask(cuda::device_buffer const& new_null_mask, size_type new_null_count, cuda::stream_ref stream = cudf::get_default_stream()); @@ -249,9 +250,10 @@ class column { * Returned by `column::release()`. */ struct contents { - std::unique_ptr data; ///< data device memory buffer - std::unique_ptr> null_mask; ///< null mask device memory buffer - std::vector> children; ///< child columns + std::unique_ptr data; ///< data device memory buffer + std::unique_ptr> + null_mask; ///< null mask device memory buffer + std::vector> children; ///< child columns }; /** @@ -326,7 +328,7 @@ class column { cudf::size_type _size{}; ///< The number of elements in the column rmm::device_buffer _data{}; ///< Dense, contiguous, type erased device memory ///< buffer containing the column elements - cuda::device_buffer _null_mask = cudf::create_null_mask( + cuda::device_buffer _null_mask = cudf::create_null_mask( 0, cudf::mask_state::UNALLOCATED); ///< Bitmask used to represent null values. ///< May be empty if `null_count() == 0` mutable cudf::size_type _null_count{}; ///< The number of null elements diff --git a/cpp/include/cudf/column/column_factories.hpp b/cpp/include/cudf/column/column_factories.hpp index f3cad6e1679f..0ef043bf2117 100644 --- a/cpp/include/cudf/column/column_factories.hpp +++ b/cpp/include/cudf/column/column_factories.hpp @@ -442,7 +442,7 @@ std::unique_ptr make_strings_column(size_type num_strings, std::unique_ptr offsets_column, rmm::device_buffer&& chars_buffer, size_type null_count, - cuda::device_buffer&& null_mask); + cuda::device_buffer&& null_mask); /** * @brief Construct a LIST type column given offsets column, child column, null mask and null @@ -497,11 +497,12 @@ std::unique_ptr make_strings_column(size_type num_strings, * @param null_mask The bits specifying the null lists in device memory. * @return Constructed lists column */ -std::unique_ptr make_lists_column(size_type num_rows, - std::unique_ptr offsets_column, - std::unique_ptr child_column, - size_type null_count, - cuda::device_buffer&& null_mask); +std::unique_ptr make_lists_column( + size_type num_rows, + std::unique_ptr offsets_column, + std::unique_ptr child_column, + size_type null_count, + cuda::device_buffer&& null_mask); /** * @brief Create an empty LIST column @@ -540,7 +541,7 @@ std::unique_ptr make_structs_column( size_type num_rows, std::vector>&& child_columns, size_type null_count, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); @@ -574,7 +575,7 @@ std::unique_ptr create_structs_hierarchy( size_type num_rows, std::vector>&& child_columns, size_type null_count, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); diff --git a/cpp/include/cudf/concatenate.hpp b/cpp/include/cudf/concatenate.hpp index 6cdc0757958a..a1b1ffcb6fc8 100644 --- a/cpp/include/cudf/concatenate.hpp +++ b/cpp/include/cudf/concatenate.hpp @@ -31,7 +31,7 @@ namespace CUDF_EXPORT cudf { /** * @brief Concatenates `views[i]`'s bitmask from the bits * `[views[i].offset(), views[i].offset() + views[i].size())` for all elements - * `views` into a `cuda::device_buffer` + * `views` into a `cuda::device_buffer` * * Returns an empty buffer if the column is not nullable. * @@ -40,7 +40,7 @@ namespace CUDF_EXPORT cudf { * @param stream CUDA stream used for device memory operations and kernel launches * @return Bitmasks of all the column views in the views vector */ -cuda::device_buffer concatenate_masks( +cuda::device_buffer concatenate_masks( std::span views, cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); diff --git a/cpp/include/cudf/detail/concatenate_masks.hpp b/cpp/include/cudf/detail/concatenate_masks.hpp index 574b2401cc47..f1ff8f56bea8 100644 --- a/cpp/include/cudf/detail/concatenate_masks.hpp +++ b/cpp/include/cudf/detail/concatenate_masks.hpp @@ -56,9 +56,9 @@ size_type concatenate_masks(host_span views, * * @param stream CUDA stream used for device memory operations and kernel launches. */ -cuda::device_buffer concatenate_masks(std::span views, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr); +cuda::device_buffer concatenate_masks(std::span views, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr); } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/detail/null_mask.cuh b/cpp/include/cudf/detail/null_mask.cuh index 7666d042a54d..cde07f06f22f 100644 --- a/cpp/include/cudf/detail/null_mask.cuh +++ b/cpp/include/cudf/detail/null_mask.cuh @@ -256,7 +256,7 @@ rmm::device_uvector inplace_segmented_bitmask_binop( * @param stream CUDA stream used for device memory operations and kernel launches */ template -std::pair, size_type> bitmask_binop( +std::pair, size_type> bitmask_binop( Binop op, host_span masks, host_span masks_begin_bits, @@ -267,21 +267,21 @@ std::pair, size_type> bitmask_binop( auto dest_mask = cudf::create_null_mask(mask_size_bits, cudf::mask_state::UNINITIALIZED, stream, mr); auto null_count = - mask_size_bits - inplace_bitmask_binop( - op, - device_span(reinterpret_cast(dest_mask.data()), - num_bitmask_words(mask_size_bits)), - masks, - masks_begin_bits, - mask_size_bits, - stream, - mr.as_temporary()); + mask_size_bits - inplace_bitmask_binop(op, + device_span( + dest_mask.data(), num_bitmask_words(mask_size_bits)), + masks, + masks_begin_bits, + mask_size_bits, + stream, + mr.as_temporary()); return std::pair(std::move(dest_mask), null_count); } template -std::pair>>, std::vector> +std::pair>>, + std::vector> segmented_bitmask_binop(Binop op, host_span masks, host_span masks_begin_bits, @@ -303,15 +303,14 @@ segmented_bitmask_binop(Binop op, auto const num_segments = segment_offsets.size() - 1; auto const temp_mr = mr.get_temporary_mr(); - std::vector>> h_destination_masks; + std::vector>> h_destination_masks; h_destination_masks.reserve(num_segments); std::vector h_destination_masks_ptrs; h_destination_masks_ptrs.reserve(num_segments); for (size_t i = 0; i < segment_offsets.size() - 1; i++) { - h_destination_masks.push_back(std::make_unique>( + h_destination_masks.push_back(std::make_unique>( cudf::create_null_mask(num_bytes * CHAR_BIT, cudf::mask_state::UNINITIALIZED, stream, mr))); - h_destination_masks_ptrs.push_back( - reinterpret_cast(h_destination_masks.back()->data())); + h_destination_masks_ptrs.push_back(h_destination_masks.back()->data()); } auto destination_masks = cudf::detail::make_device_uvector_async(h_destination_masks_ptrs, stream, temp_mr); @@ -804,7 +803,7 @@ std::vector segmented_null_count(bitmask_type const* bitmask, * @return A pair containing the reduced null mask and number of nulls. */ template -std::pair, size_type> segmented_null_mask_reduction( +std::pair, size_type> segmented_null_mask_reduction( bitmask_type const* bitmask, OffsetIterator first_bit_indices_begin, OffsetIterator first_bit_indices_end, diff --git a/cpp/include/cudf/detail/null_mask.hpp b/cpp/include/cudf/detail/null_mask.hpp index 1159cd159648..1ac8e790ab2a 100644 --- a/cpp/include/cudf/detail/null_mask.hpp +++ b/cpp/include/cudf/detail/null_mask.hpp @@ -23,10 +23,10 @@ namespace detail { * * @param stream CUDA stream used for device memory operations and kernel launches. */ -cuda::device_buffer create_null_mask(size_type size, - mask_state state, - cuda::stream_ref stream, - cudf::memory_resources mr); +cuda::device_buffer create_null_mask(size_type size, + mask_state state, + cuda::stream_ref stream, + cudf::memory_resources mr); /** * @copydoc cudf::set_null_mask(bitmask_type*, size_type, size_type, bool) @@ -162,20 +162,20 @@ std::vector segmented_null_count(bitmask_type const* bitmask, * * @param stream CUDA stream used for device memory operations and kernel launches. */ -cuda::device_buffer copy_bitmask(bitmask_type const* mask, - size_type begin_bit, - size_type end_bit, - cuda::stream_ref stream, - cudf::memory_resources mr); +cuda::device_buffer copy_bitmask(bitmask_type const* mask, + size_type begin_bit, + size_type end_bit, + cuda::stream_ref stream, + cudf::memory_resources mr); /** * @copydoc cudf::copy_bitmask(column_view const& view, cudf::memory_resources) * * @param stream CUDA stream used for device memory operations and kernel launches. */ -cuda::device_buffer copy_bitmask(column_view const& view, - cuda::stream_ref stream, - cudf::memory_resources mr); +cuda::device_buffer copy_bitmask(column_view const& view, + cuda::stream_ref stream, + cudf::memory_resources mr); /** * @copydoc bitmask_and(host_span, host_span const, @@ -183,7 +183,7 @@ cuda::device_buffer copy_bitmask(column_view const& view, * * @param stream CUDA stream used for device memory operations and kernel launches */ -std::pair, size_type> bitmask_and( +std::pair, size_type> bitmask_and( host_span masks, host_span masks_begin_bits, size_type mask_size_bits, @@ -195,14 +195,14 @@ std::pair, size_type> bitmask_and( * * @param[in] stream CUDA stream used for device memory operations and kernel launches. */ -std::pair, size_type> bitmask_and(table_view const& view, - cuda::stream_ref stream, - cudf::memory_resources mr); +std::pair, size_type> bitmask_and( + table_view const& view, cuda::stream_ref stream, cudf::memory_resources mr); /** * @copydoc cudf::segmented_bitmask_and */ -std::pair>>, std::vector> +std::pair>>, + std::vector> segmented_bitmask_and(host_span colviews, host_span segment_offsets, cuda::stream_ref stream, @@ -211,7 +211,8 @@ segmented_bitmask_and(host_span colviews, /** * @copydoc cudf::segmented_bitmask_and */ -std::pair>>, std::vector> +std::pair>>, + std::vector> segmented_bitmask_and(host_span masks, host_span segment_offsets, size_type mask_size_bits, @@ -222,9 +223,9 @@ segmented_bitmask_and(host_span masks, * * @param[in] stream CUDA stream used for device memory operations and kernel launches. */ -std::pair, size_type> bitmask_or(table_view const& view, - cuda::stream_ref stream, - cudf::memory_resources mr); +std::pair, size_type> bitmask_or(table_view const& view, + cuda::stream_ref stream, + cudf::memory_resources mr); /** * @brief Performs a bitwise AND of the specified bitmasks, diff --git a/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh b/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh index ae24f49abfba..dfde02b5f808 100644 --- a/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh +++ b/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh @@ -79,7 +79,7 @@ struct preprocessed_table { rmm::device_async_resource_ref>; preprocessed_table(table_device_view_owner&& table, - std::vector>&& null_buffers, + std::vector>&& null_buffers, std::vector>&& tmp_columns) : _t(std::move(table)), _null_buffers(std::move(null_buffers)), @@ -88,7 +88,7 @@ struct preprocessed_table { } table_device_view_owner _t; - std::vector> _null_buffers; + std::vector> _null_buffers; std::vector> _tmp_columns; }; diff --git a/cpp/include/cudf/detail/structs/utilities.hpp b/cpp/include/cudf/detail/structs/utilities.hpp index c10018935bbd..a5ba319f2f1a 100644 --- a/cpp/include/cudf/detail/structs/utilities.hpp +++ b/cpp/include/cudf/detail/structs/utilities.hpp @@ -32,7 +32,7 @@ struct temporary_nullable_data { /** * @brief Store the newly generated null masks (if any). */ - std::vector> new_null_masks; + std::vector> new_null_masks; /** * @brief Store the newly generated columns with new null masks (if any). diff --git a/cpp/include/cudf/detail/transform.hpp b/cpp/include/cudf/detail/transform.hpp index c0d2674dcb46..1cdf3c9fea0f 100644 --- a/cpp/include/cudf/detail/transform.hpp +++ b/cpp/include/cudf/detail/transform.hpp @@ -31,7 +31,7 @@ std::unique_ptr compute_column(table_view const& table, * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::pair>, cudf::size_type> bools_to_mask( +std::pair>, cudf::size_type> bools_to_mask( column_view const& input, cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** diff --git a/cpp/include/cudf/detail/valid_if.cuh b/cpp/include/cudf/detail/valid_if.cuh index 3dbc9be002c8..20f21dcfbdff 100644 --- a/cpp/include/cudf/detail/valid_if.cuh +++ b/cpp/include/cudf/detail/valid_if.cuh @@ -64,7 +64,7 @@ CUDF_KERNEL void valid_if_kernel( * * Bit `i` in the output mask will be set if `p(*(begin+i)) == true`. * - * If `distance(begin,end) == 0`, returns an empty `cuda::device_buffer`. + * If `distance(begin,end) == 0`, returns an empty `cuda::device_buffer`. * * @throws cudf::logic_error if `(begin > end)` * @@ -76,11 +76,12 @@ CUDF_KERNEL void valid_if_kernel( * @return A pair containing a `device_buffer` with the new bitmask and its null count */ template -std::pair, size_type> valid_if(InputIterator begin, - InputIterator end, - Predicate p, - cuda::stream_ref stream, - cudf::memory_resources resources) +std::pair, size_type> valid_if( + InputIterator begin, + InputIterator end, + Predicate p, + cuda::stream_ref stream, + cudf::memory_resources resources) { CUDF_EXPECTS(begin <= end, "Invalid range."); @@ -97,7 +98,7 @@ std::pair, size_type> valid_if(InputIterator begi grid_1d grid{size, block_size}; valid_if_kernel<<>>( - reinterpret_cast(null_mask.data()), begin, size, p, valid_count.data()); + null_mask.data(), begin, size, p, valid_count.data()); null_count = size - valid_count.value(stream); } diff --git a/cpp/include/cudf/dictionary/dictionary_factories.hpp b/cpp/include/cudf/dictionary/dictionary_factories.hpp index 2a5a43c421ad..19553d8c80c0 100644 --- a/cpp/include/cudf/dictionary/dictionary_factories.hpp +++ b/cpp/include/cudf/dictionary/dictionary_factories.hpp @@ -76,7 +76,7 @@ std::unique_ptr make_dictionary_column( */ std::unique_ptr make_dictionary_column(std::unique_ptr keys_column, std::unique_ptr indices_column, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, size_type null_count); /** diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index 8ee613641609..ebf62cccc88b 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -227,7 +227,7 @@ std::unique_ptr scatter(scalar const& slr, auto lv = static_cast(&slr); bool slr_valid = slr.is_valid(stream); - cuda::device_buffer null_mask = + cuda::device_buffer null_mask = slr_valid ? cudf::create_null_mask(1, mask_state::UNALLOCATED, stream, mr) : cudf::create_null_mask(1, mask_state::ALL_NULL, stream, mr); auto offset_column = @@ -240,7 +240,7 @@ std::unique_ptr scatter(scalar const& slr, auto wrapped = column_view(data_type{type_id::LIST}, 1, nullptr, - reinterpret_cast(null_mask.data()), + null_mask.data(), slr_valid ? 0 : 1, 0, {offset_column->view(), lv->view()}); diff --git a/cpp/include/cudf/null_mask.hpp b/cpp/include/cudf/null_mask.hpp index b0c4c9dbb971..d02874beef4f 100644 --- a/cpp/include/cudf/null_mask.hpp +++ b/cpp/include/cudf/null_mask.hpp @@ -55,6 +55,21 @@ size_type state_null_count(mask_state state, size_type size); std::size_t bitmask_allocation_size_bytes(size_type number_of_bits, std::size_t padding_boundary = 64); +/** + * @brief Computes the required words necessary to represent the specified + * number of bits with a given padding boundary + * + * @note The Arrow specification for the null bitmask requires a 64B padding + * boundary. + * + * @param number_of_bits The number of bits that need to be represented + * @param padding_boundary The value returned will be rounded up to a multiple + * of this value + * @return The necessary number of words + */ +std::size_t bitmask_allocation_size_words(size_type number_of_bits, + std::size_t padding_boundary = 64 / sizeof(bitmask_type)); + /** * @brief Returns the number of `bitmask_type` words required to represent the * specified number of bits @@ -79,9 +94,10 @@ size_type num_bitmask_words(size_type number_of_bits); * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Memory resources used for temporary allocations and the returned device_buffer * @return A `device_buffer` for use as a null bitmask - * satisfying the desired size and state + * satisfying the desired size and state. The buffer contains `bitmask_type` words, + * including padding to the boundary used by `bitmask_allocation_size_bytes`. */ -cuda::device_buffer create_null_mask( +cuda::device_buffer create_null_mask( size_type size, mask_state state, cuda::stream_ref stream = cudf::get_default_stream(), @@ -165,7 +181,7 @@ void set_null_masks_unsafe(cudf::host_span bitmasks, * @return A `device_buffer` containing the bits * `[begin_bit, end_bit)` from `mask`. */ -cuda::device_buffer copy_bitmask( +cuda::device_buffer copy_bitmask( bitmask_type const* mask, size_type begin_bit, size_type end_bit, @@ -184,7 +200,7 @@ cuda::device_buffer copy_bitmask( * @return A `device_buffer` containing the bits * `[view.offset(), view.offset() + view.size())` from `view`'s bitmask. */ -cuda::device_buffer copy_bitmask( +cuda::device_buffer copy_bitmask( column_view const& view, cuda::stream_ref stream = cudf::get_default_stream(), cudf::memory_resources mr = cudf::get_current_device_resource_ref()); @@ -201,7 +217,7 @@ cuda::device_buffer copy_bitmask( * @param mr Memory resources used for temporary allocations and the returned device_buffer * @return A pair of resulting bitmask and count of unset bits */ -std::pair, size_type> bitmask_and( +std::pair, size_type> bitmask_and( table_view const& view, cuda::stream_ref stream = cudf::get_default_stream(), cudf::memory_resources mr = cudf::get_current_device_resource_ref()); @@ -218,7 +234,7 @@ std::pair, size_type> bitmask_and( * @param mr Memory resources used for temporary allocations and the returned device_buffer * @return A pair of resulting bitmask of size mask_size and the count of unset bits */ -std::pair, size_type> bitmask_and( +std::pair, size_type> bitmask_and( host_span masks, host_span begin_bits, size_type mask_size, @@ -249,7 +265,8 @@ std::pair, size_type> bitmask_and( * @param mr Memory resources used for temporary allocations and the returned device_buffer * @return A pair of vectors containing resulting bitmask and count of unset bits for each segment */ -std::pair>>, std::vector> +std::pair>>, + std::vector> segmented_bitmask_and(host_span colviews, host_span segment_offsets, cuda::stream_ref stream = cudf::get_default_stream(), @@ -279,7 +296,8 @@ segmented_bitmask_and(host_span colviews, * @param mr Memory resources used for temporary allocations and the returned device_buffer * @return A pair of vectors containing resulting bitmask and count of unset bits for each segment */ -std::pair>>, std::vector> +std::pair>>, + std::vector> segmented_bitmask_and(host_span masks, host_span segment_offsets, size_type mask_size_bits, @@ -298,7 +316,7 @@ segmented_bitmask_and(host_span masks, * @param mr Memory resources used for temporary allocations and the returned device_buffer * @return A pair of resulting bitmask and count of unset bits */ -std::pair, size_type> bitmask_or( +std::pair, size_type> bitmask_or( table_view const& view, cuda::stream_ref stream = cudf::get_default_stream(), cudf::memory_resources mr = cudf::get_current_device_resource_ref()); diff --git a/cpp/include/cudf/transform.hpp b/cpp/include/cudf/transform.hpp index d62662eb944f..53f3130dbcc1 100644 --- a/cpp/include/cudf/transform.hpp +++ b/cpp/include/cudf/transform.hpp @@ -501,7 +501,7 @@ std::unique_ptr compute_table_jit( * null count obtained from input considering `true` represent `valid`/`1` and * `false` represent `invalid`/`0`. */ -std::pair>, cudf::size_type> bools_to_mask( +std::pair>, cudf::size_type> bools_to_mask( column_view const& input, cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp index 6f9d9bbf7b51..484b1a212eef 100644 --- a/cpp/include/cudf_test/column_wrapper.hpp +++ b/cpp/include/cudf_test/column_wrapper.hpp @@ -285,23 +285,24 @@ std::pair, cudf::size_type> make_null_mask_vector(Vali * @param end The end of the validity indicator sequence * @param stream CUDA stream used for device memory operations * @param mr Memory resources used to allocate the returned buffer - * @return cuda::device_buffer Contains a bitmask where bits are set for every + * @return cuda::device_buffer Contains a bitmask where bits are set for every * element in `[begin,end)` that evaluated to `true`. */ template -std::pair, cudf::size_type> make_null_mask( +std::pair, cudf::size_type> make_null_mask( ValidityIterator begin, ValidityIterator end, cuda::stream_ref stream = cudf::test::get_default_stream(), cudf::memory_resources mr = cudf::get_current_device_resource_ref()) { auto [null_mask, null_count] = make_null_mask_vector(begin, end); - auto const* data = reinterpret_cast(null_mask.data()); - cuda::device_buffer d_mask{ + auto const* data = null_mask.data(); + cuda::device_buffer d_mask{ stream, mr.get_output_mr(), data, - data + cudf::bitmask_allocation_size_bytes(cudf::distance(begin, end))}; + data + + cudf::bitmask_allocation_size_bytes(cudf::distance(begin, end)) / sizeof(cudf::bitmask_type)}; stream.sync(); // wait for async H2D before host source is destroyed return {std::move(d_mask), null_count}; } @@ -1878,7 +1879,7 @@ class lists_column_wrapper : public detail::column_wrapper { std::unique_ptr&& offsets, std::unique_ptr&& values, size_type null_count, - cuda::device_buffer&& null_mask) + cuda::device_buffer&& null_mask) { // construct the list column wrapped = make_lists_column( diff --git a/cpp/include/cudf_test/nanoarrow_utils.hpp b/cpp/include/cudf_test/nanoarrow_utils.hpp index 51855161e221..205572342e9d 100644 --- a/cpp/include/cudf_test/nanoarrow_utils.hpp +++ b/cpp/include/cudf_test/nanoarrow_utils.hpp @@ -194,11 +194,11 @@ void populate_from_col(ArrowArray* arr, ArrowArrayBuffer(arr, 1), ArrowBufferDeallocator( [](ArrowBufferAllocator* alloc, uint8_t*, int64_t) { - auto buf = - reinterpret_cast>*>(alloc->private_data); + auto buf = reinterpret_cast>*>( + alloc->private_data); delete buf; }, - new std::unique_ptr>(std::move(bitmask.first))))); + new std::unique_ptr>(std::move(bitmask.first))))); ArrowArrayBuffer(arr, 1)->size_bytes = cudf::bitmask_allocation_size_bytes(view.size()); ArrowArrayBuffer(arr, 1)->data = ptr; } diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 56d0d1672f35..63d73222a43d 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -65,7 +65,7 @@ bool is_supported_operation(data_type out, data_type lhs, data_type rhs, binary_ /** * @brief Computes output valid mask for op between a column and a scalar */ -std::pair, size_type> scalar_col_valid_mask_and( +std::pair, size_type> scalar_col_valid_mask_and( column_view const& col, scalar const& s, cuda::stream_ref stream, diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu index 30822fd8582d..4d8cea1fef4b 100644 --- a/cpp/src/binaryop/compiled/binary_ops.cu +++ b/cpp/src/binaryop/compiled/binary_ops.cu @@ -53,7 +53,7 @@ struct scalar_as_column_view { // Null scalar needs a single-element null mask (kept alive by an auxiliary column) as its // validity (bool) cannot be reinterpreted as `bitmask_type` without reading out of bounds. auto null_mask = cudf::detail::create_null_mask(1, cudf::mask_state::ALL_NULL, stream, mr); - auto const* null_mask_ptr = reinterpret_cast(null_mask.data()); + auto const* null_mask_ptr = null_mask.data(); auto aux_col = std::make_unique( data_type{type_id::INT8}, 0, rmm::device_buffer{}, std::move(null_mask), 1); auto col_v = column_view(s.type(), 1, h_scalar_type_view.data(), null_mask_ptr, 1); @@ -89,7 +89,7 @@ scalar_as_column_view::return_type scalar_as_column_view::operator()(null_mask.data()); + auto const* null_mask_ptr = null_mask.data(); auto col_v = cudf::column_view( s.type(), 1, h_scalar_type_view.data(), null_mask_ptr, 1, 0, {offsets_column->view()}); std::vector> children; diff --git a/cpp/src/bitmask/null_mask.cu b/cpp/src/bitmask/null_mask.cu index 2442a4ecf8e0..533be8621d93 100644 --- a/cpp/src/bitmask/null_mask.cu +++ b/cpp/src/bitmask/null_mask.cu @@ -49,12 +49,18 @@ size_type state_null_count(mask_state state, size_type size) // Computes required allocation size of a bitmask std::size_t bitmask_allocation_size_bytes(size_type number_of_bits, std::size_t padding_boundary) { - CUDF_EXPECTS(padding_boundary > 0, "Invalid padding boundary"); - auto necessary_bytes = cudf::util::div_rounding_up_safe(number_of_bits, CHAR_BIT); + return bitmask_allocation_size_words(number_of_bits, padding_boundary / sizeof(bitmask_type)) * + sizeof(bitmask_type); +} - auto padded_bytes = padding_boundary * cudf::util::div_rounding_up_safe( - necessary_bytes, padding_boundary); - return padded_bytes; +// Computes required alloation size of a bitmask +std::size_t bitmask_allocation_size_words(size_type number_of_bits, std::size_t padding_boundary) +{ + CUDF_EXPECTS(padding_boundary > 0, "Invalid padding boundary"); + auto necessary_words = num_bitmask_words(number_of_bits); + auto padded_words = padding_boundary * cudf::util::div_rounding_up_safe( + necessary_words, padding_boundary); + return padded_words; } // Computes number of *actual* bitmask_type elements needed @@ -67,22 +73,23 @@ size_type num_bitmask_words(size_type number_of_bits) namespace detail { // Create a device_buffer for a null mask -cuda::device_buffer create_null_mask(size_type size, - mask_state state, - cuda::stream_ref stream, - cudf::memory_resources mr) +cuda::device_buffer create_null_mask(size_type size, + mask_state state, + cuda::stream_ref stream, + cudf::memory_resources mr) { - size_type mask_size{0}; + std::size_t mask_size{0}; - if (state != mask_state::UNALLOCATED) { mask_size = bitmask_allocation_size_bytes(size); } + if (state != mask_state::UNALLOCATED) { mask_size = bitmask_allocation_size_words(size); } auto env = cuda::std::execution::prop{cuda::allocation_alignment, alignof(bitmask_type)}; - cuda::device_buffer mask(stream, mr.get_output_mr(), mask_size, cuda::no_init, env); + cuda::device_buffer mask( + stream, mr.get_output_mr(), mask_size, cuda::no_init, env); if (mask_size > 0 && state != mask_state::UNINITIALIZED) { uint8_t fill_value = (state == mask_state::ALL_VALID) ? 0xff : 0x00; - CUDF_CUDA_TRY(cudaMemsetAsync( - reinterpret_cast(mask.data()), fill_value, mask_size, stream.get())); + CUDF_CUDA_TRY( + cudaMemsetAsync(mask.data(), fill_value, mask_size * sizeof(bitmask_type), stream.get())); } return mask; @@ -276,10 +283,10 @@ void set_null_mask(bitmask_type* bitmask, } // namespace detail // Create a device_buffer for a null mask -cuda::device_buffer create_null_mask(size_type size, - mask_state state, - cuda::stream_ref stream, - cudf::memory_resources mr) +cuda::device_buffer create_null_mask(size_type size, + mask_state state, + cuda::stream_ref stream, + cudf::memory_resources mr) { CUDF_FUNC_RANGE(); return detail::create_null_mask(size, state, stream, mr); @@ -357,41 +364,36 @@ CUDF_KERNEL void copy_offset_bitmask(bitmask_type* __restrict__ destination, } // namespace // Create a bitmask from a specific range -cuda::device_buffer copy_bitmask(bitmask_type const* mask, - size_type begin_bit, - size_type end_bit, - cuda::stream_ref stream, - cudf::memory_resources mr) +cuda::device_buffer copy_bitmask(bitmask_type const* mask, + size_type begin_bit, + size_type end_bit, + cuda::stream_ref stream, + cudf::memory_resources mr) { CUDF_EXPECTS(begin_bit >= 0 and begin_bit <= end_bit, "Invalid bit range."); auto dest_mask = cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, mr); - auto num_bytes = bitmask_allocation_size_bytes(end_bit - begin_bit); - if ((mask == nullptr) || (num_bytes == 0)) { return dest_mask; } + auto num_words = bitmask_allocation_size_words(end_bit - begin_bit); + if ((mask == nullptr) || (num_words == 0)) { return dest_mask; } if (begin_bit == 0) { - auto const data = reinterpret_cast(mask); - auto env = cuda::std::execution::prop{cuda::allocation_alignment, alignof(bitmask_type)}; - dest_mask = - cuda::device_buffer{stream, mr.get_output_mr(), data, data + num_bytes, env}; + auto env = cuda::std::execution::prop{cuda::allocation_alignment, alignof(bitmask_type)}; + dest_mask = cuda::device_buffer{ + stream, mr.get_output_mr(), mask, mask + num_words, env}; } else { auto number_of_mask_words = num_bitmask_words(end_bit - begin_bit); dest_mask = cudf::create_null_mask(end_bit - begin_bit, cudf::mask_state::UNINITIALIZED, stream, mr); cudf::detail::grid_1d config(number_of_mask_words, 256); copy_offset_bitmask<<>>( - reinterpret_cast(dest_mask.data()), - mask, - begin_bit, - end_bit, - number_of_mask_words); + dest_mask.data(), mask, begin_bit, end_bit, number_of_mask_words); CUDF_CHECK_CUDA(stream.get()); } return dest_mask; } // Create a bitmask from a column view -cuda::device_buffer copy_bitmask(column_view const& view, - cuda::stream_ref stream, - cudf::memory_resources mr) +cuda::device_buffer copy_bitmask(column_view const& view, + cuda::stream_ref stream, + cudf::memory_resources mr) { auto null_mask = cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, mr); if (view.nullable()) { @@ -619,7 +621,7 @@ cudf::size_type inplace_bitmask_and(device_span dest_mask, } // Bitwise AND of the masks -std::pair, size_type> bitmask_and( +std::pair, size_type> bitmask_and( host_span masks, host_span begin_bits, size_type mask_size, @@ -636,9 +638,9 @@ std::pair, size_type> bitmask_and( } // Returns the bitwise AND of the null masks of all columns in the table view -std::pair, size_type> bitmask_and(table_view const& view, - cuda::stream_ref stream, - cudf::memory_resources mr) +std::pair, size_type> bitmask_and(table_view const& view, + cuda::stream_ref stream, + cudf::memory_resources mr) { auto null_mask = cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, mr); if (view.num_rows() == 0 or view.num_columns() == 0) { @@ -671,7 +673,8 @@ std::pair, size_type> bitmask_and(table_view cons constexpr bitmask_type all_set_mask = ~bitmask_type{0}; // Returns the bitwise AND of the null masks of all columns in the same segment of the input masks -std::pair>>, std::vector> +std::pair>>, + std::vector> segmented_bitmask_and(host_span colviews, host_span segment_offsets, cuda::stream_ref stream, @@ -701,7 +704,8 @@ segmented_bitmask_and(host_span colviews, mr); } -std::pair>>, std::vector> +std::pair>>, + std::vector> segmented_bitmask_and(host_span masks, host_span segment_offsets, size_type mask_size_bits, @@ -724,9 +728,9 @@ segmented_bitmask_and(host_span masks, } // Returns the bitwise OR of the null masks of all columns in the table view -std::pair, size_type> bitmask_or(table_view const& view, - cuda::stream_ref stream, - cudf::memory_resources mr) +std::pair, size_type> bitmask_or(table_view const& view, + cuda::stream_ref stream, + cudf::memory_resources mr) { auto null_mask = cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, mr); if (view.num_rows() == 0 or view.num_columns() == 0) { @@ -836,34 +840,34 @@ size_type index_of_first_set_bit(bitmask_type const* bitmask, } // namespace detail // Create a bitmask from a specific range -cuda::device_buffer copy_bitmask(bitmask_type const* mask, - size_type begin_bit, - size_type end_bit, - cuda::stream_ref stream, - cudf::memory_resources mr) +cuda::device_buffer copy_bitmask(bitmask_type const* mask, + size_type begin_bit, + size_type end_bit, + cuda::stream_ref stream, + cudf::memory_resources mr) { CUDF_FUNC_RANGE(); return detail::copy_bitmask(mask, begin_bit, end_bit, stream, mr); } // Create a bitmask from a column view -cuda::device_buffer copy_bitmask(column_view const& view, - cuda::stream_ref stream, - cudf::memory_resources mr) +cuda::device_buffer copy_bitmask(column_view const& view, + cuda::stream_ref stream, + cudf::memory_resources mr) { CUDF_FUNC_RANGE(); return detail::copy_bitmask(view, stream, mr); } -std::pair, size_type> bitmask_and(table_view const& view, - cuda::stream_ref stream, - cudf::memory_resources mr) +std::pair, size_type> bitmask_and(table_view const& view, + cuda::stream_ref stream, + cudf::memory_resources mr) { CUDF_FUNC_RANGE(); return detail::bitmask_and(view, stream, mr); } -std::pair, size_type> bitmask_and( +std::pair, size_type> bitmask_and( host_span masks, host_span begin_bits, size_type mask_size, @@ -874,7 +878,8 @@ std::pair, size_type> bitmask_and( return detail::bitmask_and(masks, begin_bits, mask_size, stream, mr); } -std::pair>>, std::vector> +std::pair>>, + std::vector> segmented_bitmask_and(host_span colviews, host_span segment_offsets, cuda::stream_ref stream, @@ -884,7 +889,8 @@ segmented_bitmask_and(host_span colviews, return detail::segmented_bitmask_and(colviews, segment_offsets, stream, mr); } -std::pair>>, std::vector> +std::pair>>, + std::vector> segmented_bitmask_and(host_span masks, host_span segment_offsets, size_type mask_size_bits, @@ -895,9 +901,9 @@ segmented_bitmask_and(host_span masks, return detail::segmented_bitmask_and(masks, segment_offsets, mask_size_bits, stream, mr); } -std::pair, size_type> bitmask_or(table_view const& view, - cuda::stream_ref stream, - cudf::memory_resources mr) +std::pair, size_type> bitmask_or(table_view const& view, + cuda::stream_ref stream, + cudf::memory_resources mr) { CUDF_FUNC_RANGE(); return detail::bitmask_or(view, stream, mr); diff --git a/cpp/src/column/column.cu b/cpp/src/column/column.cu index 7a2a65c4a995..4e2a756128d7 100644 --- a/cpp/src/column/column.cu +++ b/cpp/src/column/column.cu @@ -41,8 +41,7 @@ column::column(column const& other, cuda::stream_ref stream, rmm::device_async_r : _type{other._type}, _size{other._size}, _data{other._data, stream, mr}, - _null_mask{cudf::detail::copy_bitmask( - reinterpret_cast(other._null_mask.data()), 0, other.size(), stream, mr)}, + _null_mask{cudf::detail::copy_bitmask(other._null_mask.data(), 0, other.size(), stream, mr)}, _null_count{other._null_count} { _children.reserve(other.num_children()); @@ -71,14 +70,15 @@ column::contents column::release() noexcept _size = 0; _null_count = 0; _type = data_type{type_id::EMPTY}; - return column::contents{std::make_unique(std::move(_data)), - std::make_unique>(std::move(_null_mask)), - std::move(_children)}; + return column::contents{ + std::make_unique(std::move(_data)), + std::make_unique>(std::move(_null_mask)), + std::move(_children)}; } std::size_t column::alloc_size() const { - return _data.size() + _null_mask.size() + + return _data.size() + _null_mask.size() * sizeof(bitmask_type) + std::transform_reduce( _children.begin(), _children.end(), std::size_t{0}, std::plus{}, [](auto const& c) { return c->alloc_size(); @@ -95,13 +95,7 @@ column_view column::view() const child_views.emplace_back(*c); } - return column_view{type(), - size(), - _data.data(), - reinterpret_cast(_null_mask.data()), - null_count(), - 0, - child_views}; + return column_view{type(), size(), _data.data(), _null_mask.data(), null_count(), 0, child_views}; } // Create mutable view @@ -116,19 +110,16 @@ mutable_column_view column::mutable_view() child_views.emplace_back(*c); } - return mutable_column_view{type(), - size(), - _data.data(), - reinterpret_cast(_null_mask.data()), - _null_count, - 0, - child_views}; + return mutable_column_view{ + type(), size(), _data.data(), _null_mask.data(), _null_count, 0, child_views}; } -void column::set_null_mask(cuda::device_buffer&& new_null_mask, size_type new_null_count) +void column::set_null_mask(cuda::device_buffer&& new_null_mask, + size_type new_null_count) { if (new_null_count > 0) { - CUDF_EXPECTS(new_null_mask.size() >= cudf::bitmask_allocation_size_bytes(this->size()), + CUDF_EXPECTS(new_null_mask.size() * sizeof(bitmask_type) >= + cudf::bitmask_allocation_size_bytes(this->size()), "Column with null values must be nullable and the null mask \ buffer size should match the size of the column."); } @@ -136,16 +127,17 @@ void column::set_null_mask(cuda::device_buffer&& new_null_mask, size_ _null_count = new_null_count; } -void column::set_null_mask(cuda::device_buffer const& new_null_mask, +void column::set_null_mask(cuda::device_buffer const& new_null_mask, size_type new_null_count, cuda::stream_ref stream) { if (new_null_count > 0) { - CUDF_EXPECTS(new_null_mask.size() >= cudf::bitmask_allocation_size_bytes(this->size()), + CUDF_EXPECTS(new_null_mask.size() * sizeof(bitmask_type) >= + cudf::bitmask_allocation_size_bytes(this->size()), "Column with null values must be nullable and the null mask \ buffer size should match the size of the column."); } - _null_mask = cuda::device_buffer{ + _null_mask = cuda::device_buffer{ stream, cudf::get_current_device_resource_ref(), new_null_mask}; // copy _null_count = new_null_count; } diff --git a/cpp/src/copying/concatenate.cu b/cpp/src/copying/concatenate.cu index a2a6531125c3..dc922d4c34f6 100644 --- a/cpp/src/copying/concatenate.cu +++ b/cpp/src/copying/concatenate.cu @@ -578,9 +578,9 @@ std::unique_ptr
concatenate(std::span tables_to_concat, return std::make_unique
(std::move(concat_columns)); } -cuda::device_buffer concatenate_masks(std::span views, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) +cuda::device_buffer concatenate_masks(std::span views, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) { bool const has_nulls = std::any_of(views.begin(), views.end(), [](column_view const col) { return col.has_nulls(); }); @@ -590,12 +590,11 @@ cuda::device_buffer concatenate_masks(std::span vi return accumulator + v.size(); }); - cuda::device_buffer null_mask = + cuda::device_buffer null_mask = cudf::detail::create_null_mask(total_element_count, mask_state::UNINITIALIZED, stream, mr); - detail::concatenate_masks(host_span{views.data(), views.size()}, - reinterpret_cast(null_mask.data()), - stream); + detail::concatenate_masks( + host_span{views.data(), views.size()}, null_mask.data(), stream); return null_mask; } @@ -605,9 +604,9 @@ cuda::device_buffer concatenate_masks(std::span vi } // namespace detail -cuda::device_buffer concatenate_masks(std::span views, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) +cuda::device_buffer concatenate_masks(std::span views, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); return detail::concatenate_masks(views, stream, mr); diff --git a/cpp/src/copying/copy_range.cu b/cpp/src/copying/copy_range.cu index a42e9afcc661..7ae0cfcd5c7e 100644 --- a/cpp/src/copying/copy_range.cu +++ b/cpp/src/copying/copy_range.cu @@ -166,12 +166,11 @@ std::unique_ptr out_of_place_copy_range_dispatch::operator()release(); auto target_indices(std::move(target_contents.children.front())); - cudf::mutable_column_view new_indices( - target_indices->type(), - dict_target.size(), - target_indices->mutable_view().head(), - reinterpret_cast(target_contents.null_mask->data()), - dict_target.null_count()); + cudf::mutable_column_view new_indices(target_indices->type(), + dict_target.size(), + target_indices->mutable_view().head(), + target_contents.null_mask->data(), + dict_target.null_count()); // TODO: route through the caller's memory resources once copy_range is migrated auto const temp_mr = cudf::get_current_device_resource_ref(); cudf::type_dispatcher(new_indices.type(), diff --git a/cpp/src/copying/scatter.cu b/cpp/src/copying/scatter.cu index abd8dae4c5d3..8dd002505956 100644 --- a/cpp/src/copying/scatter.cu +++ b/cpp/src/copying/scatter.cu @@ -259,7 +259,7 @@ struct column_scalar_scatterer_impl { std::vector> fields(fields_iter_begin, fields_iter_begin + n_fields); // Compute null mask - cuda::device_buffer null_mask = + cuda::device_buffer null_mask = target.nullable() ? detail::copy_bitmask(target, stream, mr) : detail::create_null_mask(target.size(), mask_state::UNALLOCATED, stream, mr); diff --git a/cpp/src/copying/shift.cu b/cpp/src/copying/shift.cu index f5e6c7e82aa9..5581946a3a30 100644 --- a/cpp/src/copying/shift.cu +++ b/cpp/src/copying/shift.cu @@ -37,7 +37,7 @@ inline bool __device__ out_of_bounds(size_type size, size_type idx) return idx < 0 || idx >= size; } -std::pair, size_type> create_null_mask( +std::pair, size_type> create_null_mask( column_device_view const& input, size_type offset, scalar const& fill_value, diff --git a/cpp/src/dictionary/dictionary_factories.cu b/cpp/src/dictionary/dictionary_factories.cu index f27533181c31..15287d811e90 100644 --- a/cpp/src/dictionary/dictionary_factories.cu +++ b/cpp/src/dictionary/dictionary_factories.cu @@ -68,7 +68,7 @@ std::unique_ptr make_dictionary_column(column_view const& keys_column, std::unique_ptr make_dictionary_column(std::unique_ptr keys_column, std::unique_ptr indices_column, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, size_type null_count) { CUDF_EXPECTS(!keys_column->has_nulls(), "keys column must not have nulls", std::invalid_argument); diff --git a/cpp/src/dictionary/remove_keys.cu b/cpp/src/dictionary/remove_keys.cu index 249bd37d3cc0..94bcf6dcc094 100644 --- a/cpp/src/dictionary/remove_keys.cu +++ b/cpp/src/dictionary/remove_keys.cu @@ -137,7 +137,7 @@ std::unique_ptr remove_keys_fn(dictionary_column_view const& dictionary_ }, stream, mr); - cuda::device_buffer new_null_mask = + cuda::device_buffer new_null_mask = (new_nulls.second > 0) ? std::move(new_nulls.first) : cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, mr); diff --git a/cpp/src/dictionary/set_keys.cu b/cpp/src/dictionary/set_keys.cu index dc9a8c47cacc..570df6d1901d 100644 --- a/cpp/src/dictionary/set_keys.cu +++ b/cpp/src/dictionary/set_keys.cu @@ -67,7 +67,7 @@ struct apply_indices_map_fn { struct remap_result { std::unique_ptr indices; - cuda::device_buffer null_mask = + cuda::device_buffer null_mask = cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); cudf::size_type null_count; }; diff --git a/cpp/src/groupby/common/null_mask.cu b/cpp/src/groupby/common/null_mask.cu index 9254339fa8c3..50bbb8ae1ccd 100644 --- a/cpp/src/groupby/common/null_mask.cu +++ b/cpp/src/groupby/common/null_mask.cu @@ -8,13 +8,13 @@ #include namespace cudf::groupby::detail { -std::pair, size_type> make_mask_from_validity( +std::pair, size_type> make_mask_from_validity( bool* begin, bool* end, cuda::stream_ref stream, cudf::memory_resources mr) { return cudf::detail::valid_if(begin, end, cuda::std::identity{}, stream, mr); } -std::pair, size_type> make_mask_from_counts( +std::pair, size_type> make_mask_from_counts( size_type const* begin, size_type const* end, cuda::stream_ref stream, cudf::memory_resources mr) { return cudf::detail::valid_if( diff --git a/cpp/src/groupby/common/utils.cpp b/cpp/src/groupby/common/utils.cpp index a747e3db87d7..b8445371f2d2 100644 --- a/cpp/src/groupby/common/utils.cpp +++ b/cpp/src/groupby/common/utils.cpp @@ -11,7 +11,7 @@ namespace cudf::groupby::detail { -std::pair, bitmask_type const*> compute_row_bitmask( +std::pair, bitmask_type const*> compute_row_bitmask( table_view const& keys, cuda::stream_ref stream) { auto const mr = cudf::get_current_device_resource_ref(); @@ -26,14 +26,15 @@ std::pair, bitmask_type const*> compute_row_bitma col.null_mask()}; } auto buf = cudf::copy_bitmask(col, stream, mr); - auto ptr = reinterpret_cast(buf.data()); + auto ptr = buf.data(); return {std::move(buf), ptr}; } auto [buf, null_count] = cudf::bitmask_and(keys, stream, mr); if (null_count == 0) { return {cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, mr), nullptr}; } - return {std::move(buf), reinterpret_cast(buf.data())}; + auto const ptr = buf.data(); + return {std::move(buf), ptr}; } } // namespace cudf::groupby::detail diff --git a/cpp/src/groupby/common/utils.hpp b/cpp/src/groupby/common/utils.hpp index f8c7173337dd..fe2acfededa0 100644 --- a/cpp/src/groupby/common/utils.hpp +++ b/cpp/src/groupby/common/utils.hpp @@ -24,9 +24,9 @@ namespace cudf::groupby::detail { // Non-template bridges share mask-kernel instantiations across groupby reductions. -std::pair, size_type> make_mask_from_validity( +std::pair, size_type> make_mask_from_validity( bool* begin, bool* end, cuda::stream_ref stream, cudf::memory_resources mr); -std::pair, size_type> make_mask_from_counts( +std::pair, size_type> make_mask_from_counts( size_type const* begin, size_type const* end, cuda::stream_ref stream, cudf::memory_resources mr); template @@ -64,7 +64,7 @@ inline std::vector extract_results(std::span, bitmask_type const*> compute_row_bitmask( +std::pair, bitmask_type const*> compute_row_bitmask( table_view const& keys, cuda::stream_ref stream); /// Whether the given aggregation kind is supported by hash-based groupby. diff --git a/cpp/src/groupby/hash/compute_groupby.cu b/cpp/src/groupby/hash/compute_groupby.cu index 27a26a53fec8..88629d46ec62 100644 --- a/cpp/src/groupby/hash/compute_groupby.cu +++ b/cpp/src/groupby/hash/compute_groupby.cu @@ -446,7 +446,7 @@ std::unique_ptr
compute_groupby(table_view const& keys, [[maybe_unused]] auto [row_bitmask_data, row_bitmask] = skip_rows_with_nulls ? cudf::groupby::detail::compute_row_bitmask(keys, stream) - : std::pair, bitmask_type const*>{ + : std::pair, bitmask_type const*>{ cudf::create_null_mask(0, mask_state::UNALLOCATED, stream, temp_mr), nullptr}; // Determine which grouping outputs the requested aggregations need. diff --git a/cpp/src/groupby/hash/hash_compound_agg_finalizer.cu b/cpp/src/groupby/hash/hash_compound_agg_finalizer.cu index 305a9c75718a..4a6eedc274b9 100644 --- a/cpp/src/groupby/hash/hash_compound_agg_finalizer.cu +++ b/cpp/src/groupby/hash/hash_compound_agg_finalizer.cu @@ -125,7 +125,7 @@ void hash_compound_agg_finalizer::operator()(aggregation cons count_result.begin(), count_result.end(), stream, mr.as_temporary()); if (null_count > 0) { result->set_null_mask( - cuda::device_buffer{ + cuda::device_buffer{ stream, mr.get_output_mr(), null_mask.begin(), null_mask.end()}, null_count); } diff --git a/cpp/src/groupby/hash/output_utils.cu b/cpp/src/groupby/hash/output_utils.cu index bddb2dcbfd5e..c74c1228e7cb 100644 --- a/cpp/src/groupby/hash/output_utils.cu +++ b/cpp/src/groupby/hash/output_utils.cu @@ -87,7 +87,8 @@ struct result_column_creator { return children; }; - auto [null_mask, null_count] = [&]() -> std::pair, size_type> { + auto [null_mask, + null_count] = [&]() -> std::pair, size_type> { if (output_size > 0 && nullable) { return {create_null_mask(output_size, mask_state::ALL_NULL, stream, mr), output_size}; } diff --git a/cpp/src/groupby/hash/single_pass_reductions.cuh b/cpp/src/groupby/hash/single_pass_reductions.cuh index 4b9209e0f34f..aaf0ab25d432 100644 --- a/cpp/src/groupby/hash/single_pass_reductions.cuh +++ b/cpp/src/groupby/hash/single_pass_reductions.cuh @@ -257,7 +257,7 @@ struct grouped_reduction_fn { ctx.nullable ? mask_state::ALL_VALID : mask_state::UNALLOCATED, stream, mr.get_output_mr()); - auto const mask = reinterpret_cast(null_mask.data()); + auto const mask = null_mask.data(); reduce(cuda::tabulate_output_iterator{ [out, mask, values = ctx.d_values] __device__(cuda::std::ptrdiff_t group, size_type row) { out[group] = row; @@ -294,7 +294,7 @@ struct grouped_reduction_fn { if (ctx.nullable) { null_mask = cudf::create_null_mask(ctx.num_groups, mask_state::ALL_VALID, stream, mr.get_output_mr()); - auto const mask = reinterpret_cast(null_mask.data()); + auto const mask = null_mask.data(); auto const values = cudf::detail::make_counting_transform_iterator( 0, grouped_sum_overflow_fn{ diff --git a/cpp/src/groupby/sort/group_sum_overflow.cu b/cpp/src/groupby/sort/group_sum_overflow.cu index 89a039890f90..1b8defe93676 100644 --- a/cpp/src/groupby/sort/group_sum_overflow.cu +++ b/cpp/src/groupby/sort/group_sum_overflow.cu @@ -80,7 +80,8 @@ struct group_sum_overflow_fn { // A group's struct entry is null only when every row in the group is null (mirrors group_sum): // reduce per-row validity with logical-or, then build the mask from the per-group result. - auto [null_mask, null_count] = [&]() -> std::pair, size_type> { + auto [null_mask, + null_count] = [&]() -> std::pair, size_type> { if (!values.has_nulls()) { return {cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), size_type{0}}; } diff --git a/cpp/src/groupby/streaming_groupby/common.cuh b/cpp/src/groupby/streaming_groupby/common.cuh index 954cb8c21123..4e500fbaeacb 100644 --- a/cpp/src/groupby/streaming_groupby/common.cuh +++ b/cpp/src/groupby/streaming_groupby/common.cuh @@ -369,7 +369,7 @@ struct streaming_groupby::impl { struct batch_insert_result { cuda::device_buffer target_indices; size_type new_insertions; - cuda::device_buffer bitmask_buffer; + cuda::device_buffer bitmask_buffer; }; batch_insert_result probe_and_insert(table_view const& batch_keys, cuda::stream_ref stream); diff --git a/cpp/src/groupby/streaming_groupby/insert.cuh b/cpp/src/groupby/streaming_groupby/insert.cuh index 71fcc2e74eb6..a51cbd8df474 100644 --- a/cpp/src/groupby/streaming_groupby/insert.cuh +++ b/cpp/src/groupby/streaming_groupby/insert.cuh @@ -52,7 +52,7 @@ streaming_groupby::impl::batch_insert_result streaming_groupby::impl::probe_and_ auto [bitmask_buffer, batch_bitmask] = skip_rows_with_nulls ? detail::compute_row_bitmask(batch_keys, stream) - : std::pair, bitmask_type const*>{ + : std::pair, bitmask_type const*>{ cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream), nullptr}; // Precompute batch hash values. Caching is faster than inlining the row hasher diff --git a/cpp/src/interop/from_arrow_host.cu b/cpp/src/interop/from_arrow_host.cu index 1ecd32b5bff4..0b8709c8a412 100644 --- a/cpp/src/interop/from_arrow_host.cu +++ b/cpp/src/interop/from_arrow_host.cu @@ -94,18 +94,18 @@ CUDF_KERNEL void copy_shifted_bitmask(bitmask_type* __restrict__ destination, } // copies the bitmask to device and automatically applies the offset -std::pair>, size_type> get_mask_buffer( +std::pair>, size_type> get_mask_buffer( ArrowArray const* input, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input->length == 0) { - return {std::make_unique>( + return {std::make_unique>( cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, mr)), 0}; } auto bitmap = static_cast(input->buffers[validity_buffer_idx]); if (bitmap == nullptr || input->null_count == 0) { - return {std::make_unique>( + return {std::make_unique>( cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, mr)), 0}; } @@ -135,9 +135,8 @@ std::pair>, size_type> get_mask_b auto const null_count = mask_words > 0 ? cudf::detail::count_unset_bits(mask.data(), 0, num_rows, stream) : 0; - auto const mask_bytes = reinterpret_cast(mask.data()); - return {std::make_unique>(cudf::detail::copy_bitmask( - reinterpret_cast(mask_bytes), 0, num_rows, stream, mr)), + return {std::make_unique>( + cudf::detail::copy_bitmask(mask.data(), 0, num_rows, stream, mr)), null_count}; } @@ -239,7 +238,7 @@ std::unique_ptr dispatch_copy_from_arrow_host::operator()length == 0) { return make_empty_column(type_id::STRING); } auto [mask, null_count] = !skip_mask ? get_mask_buffer(input, stream, mr) - : std::pair{std::make_unique>( + : std::pair{std::make_unique>( cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, mr)), 0}; return string_column_from_arrow_host(schema, input, std::move(mask), null_count, stream, mr); @@ -296,7 +295,7 @@ std::unique_ptr dispatch_copy_from_arrow_host::operator()>( + : std::pair{std::make_unique>( cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, mr)), 0}; @@ -363,7 +362,7 @@ std::unique_ptr dispatch_copy_from_arrow_host::operator()>( + : std::pair{std::make_unique>( cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, mr)), 0}; diff --git a/cpp/src/interop/from_arrow_host.hpp b/cpp/src/interop/from_arrow_host.hpp index 855174496f3b..ec924ba35dd2 100644 --- a/cpp/src/interop/from_arrow_host.hpp +++ b/cpp/src/interop/from_arrow_host.hpp @@ -32,7 +32,7 @@ namespace detail { std::unique_ptr string_column_from_arrow_host( ArrowSchemaView const* schema, ArrowArray const* input, - std::unique_ptr>&& mask, + std::unique_ptr>&& mask, size_type null_count, cuda::stream_ref stream, rmm::device_async_resource_ref mr); diff --git a/cpp/src/interop/from_arrow_host_strings.cu b/cpp/src/interop/from_arrow_host_strings.cu index 738ba4ffb99f..084e9674b785 100644 --- a/cpp/src/interop/from_arrow_host_strings.cu +++ b/cpp/src/interop/from_arrow_host_strings.cu @@ -44,12 +44,13 @@ namespace { constexpr int chars_buffer_idx = 2; -std::unique_ptr from_arrow_string(ArrowSchemaView const* schema, - ArrowArray const* input, - std::unique_ptr>&& mask, - size_type null_count, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) +std::unique_ptr from_arrow_string( + ArrowSchemaView const* schema, + ArrowArray const* input, + std::unique_ptr>&& mask, + size_type null_count, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) { auto [offsets_column, offset, char_data_length] = get_offsets_column(schema, input, stream, mr); @@ -69,7 +70,7 @@ constexpr int stringview_vector_idx = 1; std::unique_ptr from_arrow_stringview( ArrowSchemaView const* schema, ArrowArray const* input, - std::unique_ptr>&& mask, + std::unique_ptr>&& mask, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { @@ -104,7 +105,7 @@ std::unique_ptr from_arrow_stringview( auto d_variadic_ptrs = cudf::detail::make_device_buffer_async( variadic_ptrs, stream, cudf::get_current_device_resource_ref()); auto d_ptrs = d_variadic_ptrs.data(); - auto d_mask = reinterpret_cast(mask->data()); + auto d_mask = mask->data(); using string_index_pair = cudf::strings::detail::string_index_pair; @@ -136,7 +137,7 @@ std::unique_ptr from_arrow_stringview( std::unique_ptr string_column_from_arrow_host( ArrowSchemaView const* schema, ArrowArray const* input, - std::unique_ptr>&& mask, + std::unique_ptr>&& mask, size_type null_count, cuda::stream_ref stream, rmm::device_async_resource_ref mr) diff --git a/cpp/src/interop/to_arrow_device.cu b/cpp/src/interop/to_arrow_device.cu index 61bf7ca53b38..e7483ea277ea 100644 --- a/cpp/src/interop/to_arrow_device.cu +++ b/cpp/src/interop/to_arrow_device.cu @@ -65,6 +65,8 @@ int set_buffer(std::unique_ptr device_buf, int64_t i, ArrowArray* out) return sizeof(typename T::value_type); } else if constexpr (is_device_uvector::value) { return sizeof(typename T::value_type) * device_buf->size(); + } else if constexpr (std::is_same_v>) { + return sizeof(bitmask_type) * device_buf->size(); } else { return device_buf->size(); } diff --git a/cpp/src/interop/to_arrow_host.cu b/cpp/src/interop/to_arrow_host.cu index 60c74fadd449..f0cbbf2a3a55 100644 --- a/cpp/src/interop/to_arrow_host.cu +++ b/cpp/src/interop/to_arrow_host.cu @@ -88,8 +88,7 @@ struct dispatch_to_arrow_host { enable_hugepage(&bitmap->buffer); CUDF_CUDA_TRY(cudf::detail::memcpy_async( bitmap->buffer.data, - (column.offset() > 0) ? reinterpret_cast( - cudf::detail::copy_bitmask(column, stream, mr).data()) + (column.offset() > 0) ? cudf::detail::copy_bitmask(column, stream, mr).data() : column.null_mask(), bitmap->buffer.size_bytes, stream)); @@ -149,7 +148,7 @@ int dispatch_to_arrow_host::operator()(ArrowArray* out) const auto bitmask = detail::bools_to_mask(column, stream, mr); NANOARROW_RETURN_NOT_OK(populate_data_buffer( device_span(reinterpret_cast(bitmask.first->data()), - bitmask.first->size()), + bitmask.first->size() * sizeof(bitmask_type)), ArrowArrayBuffer(tmp.get(), fixed_width_data_buffer_idx))); ArrowArrayMove(tmp.get(), out); diff --git a/cpp/src/io/csv/durations.cu b/cpp/src/io/csv/durations.cu index cc82f9ec25b1..7387e416234a 100644 --- a/cpp/src/io/csv/durations.cu +++ b/cpp/src/io/csv/durations.cu @@ -171,7 +171,8 @@ struct dispatch_from_durations_fn { auto d_column = *column; // copy null mask - cuda::device_buffer null_mask = cudf::detail::copy_bitmask(durations, stream, mr); + cuda::device_buffer null_mask = + cudf::detail::copy_bitmask(durations, stream, mr); // build offsets column auto offsets_transformer_itr = diff --git a/cpp/src/io/json/host_tree_algorithms.cu b/cpp/src/io/json/host_tree_algorithms.cu index 8e2b5d2843ae..58314a0193f8 100644 --- a/cpp/src/io/json/host_tree_algorithms.cu +++ b/cpp/src/io/json/host_tree_algorithms.cu @@ -1014,7 +1014,7 @@ void scatter_offsets(tree_meta_t const& tree, columns_data[col_id] = json_column_data{col.string_offsets.data(), col.string_lengths.data(), col.child_offsets.data(), - reinterpret_cast(col.validity.data())}; + col.validity.data()}; } auto d_ignore_vals = cudf::detail::make_device_uvector_async( diff --git a/cpp/src/io/json/json_column.cu b/cpp/src/io/json/json_column.cu index 9030fe0ae8f3..8314dfbf4179 100644 --- a/cpp/src/io/json/json_column.cu +++ b/cpp/src/io/json/json_column.cu @@ -293,15 +293,15 @@ std::pair, std::vector> device_json_co { CUDF_FUNC_RANGE(); auto validity_size_check = [](device_json_column& json_col) { - CUDF_EXPECTS(json_col.validity.size() >= bitmask_allocation_size_bytes(json_col.num_rows), + CUDF_EXPECTS(json_col.validity.size() * sizeof(bitmask_type) >= + bitmask_allocation_size_bytes(json_col.num_rows), "valid_count is too small"); }; - auto make_validity = - [stream, validity_size_check]( - device_json_column& json_col) -> std::pair, size_type> { + auto make_validity = [stream, validity_size_check](device_json_column& json_col) + -> std::pair, size_type> { validity_size_check(json_col); - auto null_count = cudf::detail::null_count( - reinterpret_cast(json_col.validity.data()), 0, json_col.num_rows, stream); + auto null_count = + cudf::detail::null_count(json_col.validity.data(), 0, json_col.num_rows, stream); // full null_mask is always required for parse_data return {std::move(json_col.validity), null_count}; // Note: json_col modified here, moves this memory diff --git a/cpp/src/io/json/nested_json.hpp b/cpp/src/io/json/nested_json.hpp index 36766df4a1d4..8724a0a82368 100644 --- a/cpp/src/io/json/nested_json.hpp +++ b/cpp/src/io/json/nested_json.hpp @@ -154,7 +154,7 @@ struct device_json_column { rmm::device_uvector child_offsets; // Validity bitmap - cuda::device_buffer validity = + cuda::device_buffer validity = cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); // Map of child columns, if applicable. diff --git a/cpp/src/io/json/nested_json_gpu.cu b/cpp/src/io/json/nested_json_gpu.cu index 46886f93e2b4..22edac574007 100644 --- a/cpp/src/io/json/nested_json_gpu.cu +++ b/cpp/src/io/json/nested_json_gpu.cu @@ -2100,17 +2100,15 @@ std::pair, std::vector> json_column_to // Range of orchestrating/encapsulating function CUDF_FUNC_RANGE(); - auto make_validity = - [stream, - mr](json_column const& json_col) -> std::pair, size_type> { + auto make_validity = [stream, mr](json_column const& json_col) + -> std::pair, size_type> { auto const null_count = json_col.current_offset - json_col.valid_count; if (null_count == 0) { return {cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), null_count}; } - auto const data = reinterpret_cast(json_col.validity.data()); - return {cudf::detail::copy_bitmask( - reinterpret_cast(data), 0, json_col.current_offset, stream, mr), - null_count}; + return { + cudf::detail::copy_bitmask(json_col.validity.data(), 0, json_col.current_offset, stream, mr), + null_count}; }; auto get_child_schema = [schema](auto child_name) -> std::optional { diff --git a/cpp/src/io/json/write_json.cu b/cpp/src/io/json/write_json.cu index 0ce6c93066fb..1bc6e6133fa0 100644 --- a/cpp/src/io/json/write_json.cu +++ b/cpp/src/io/json/write_json.cu @@ -156,7 +156,7 @@ struct escape_strings_fn { std::unique_ptr make_strings_column(size_type size, size_type null_count, - cuda::device_buffer null_mask, + cuda::device_buffer null_mask, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { @@ -176,7 +176,7 @@ struct escape_strings_fn { std::unique_ptr make_escaped_json_strings(column_device_view const& d_column, size_type size, size_type null_count, - cuda::device_buffer null_mask, + cuda::device_buffer null_mask, bool append_colon, bool escaped_utf8, cuda::stream_ref stream, diff --git a/cpp/src/io/json/write_json.hpp b/cpp/src/io/json/write_json.hpp index 83e75be7faea..3f9dcab44600 100644 --- a/cpp/src/io/json/write_json.hpp +++ b/cpp/src/io/json/write_json.hpp @@ -27,7 +27,7 @@ namespace io::json::detail { std::unique_ptr make_escaped_json_strings(column_device_view const& d_column, size_type size, size_type null_count, - cuda::device_buffer null_mask, + cuda::device_buffer null_mask, bool append_colon, bool escaped_utf8, cuda::stream_ref stream, diff --git a/cpp/src/io/orc/reader_impl_decode.cu b/cpp/src/io/orc/reader_impl_decode.cu index ac168729f97c..2e28db94afc9 100644 --- a/cpp/src/io/orc/reader_impl_decode.cu +++ b/cpp/src/io/orc/reader_impl_decode.cu @@ -318,7 +318,7 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks, auto merged_null_mask = cudf::detail::create_null_mask( parent_mask_len, mask_state::ALL_NULL, cuda::stream_ref(stream), mr); - auto merged_mask = reinterpret_cast(merged_null_mask.data()); + auto merged_mask = merged_null_mask.data(); uint32_t* dst_idx_ptr = dst_idx.data(); // Copy child valid bits from child column to valid indexes, this will merge both child // and parent null masks diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 9464f03b4e3a..bac431bd51b9 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -1147,7 +1147,7 @@ struct cast_variant_fn { size_type num_rows; data_type desired_type; bitmask_type* d_null_mask; - cuda::device_buffer null_mask{ + cuda::device_buffer null_mask{ cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)}; cuda::stream_ref stream; rmm::device_async_resource_ref mr; @@ -1408,7 +1408,7 @@ std::unique_ptr get_variant_field(column_view const& variant_column, variant_column.nullable() ? cudf::detail::copy_bitmask(variant_column, stream, mr) : cudf::create_null_mask(variant_column.size(), mask_state::ALL_VALID, stream, mr); - auto* d_null_mask = reinterpret_cast(null_mask.data()); + auto* d_null_mask = null_mask.data(); auto grid = cudf::detail::grid_1d{num_rows, block_size}; @@ -1513,7 +1513,7 @@ std::unique_ptr cast_variant(column_view const& values, auto null_mask = values.nullable() ? cudf::detail::copy_bitmask(values, stream, mr) : cudf::create_null_mask(num_rows, mask_state::ALL_VALID, stream, mr); - auto* d_null_mask = reinterpret_cast(null_mask.data()); + auto* d_null_mask = null_mask.data(); return cudf::type_dispatcher( desired_type, @@ -1542,7 +1542,7 @@ std::unique_ptr get_variant_type_id(column_view const& values, auto null_mask = values.nullable() ? cudf::detail::copy_bitmask(values, stream, mr) : cudf::create_null_mask(num_rows, mask_state::ALL_VALID, stream, mr); - auto* d_null_mask = reinterpret_cast(null_mask.data()); + auto* d_null_mask = null_mask.data(); rmm::device_buffer data{static_cast(num_rows) * sizeof(uint8_t), stream, mr}; diff --git a/cpp/src/io/parquet/stats_filter_helpers.hpp b/cpp/src/io/parquet/stats_filter_helpers.hpp index 2c63c33c4168..3a02c4224458 100644 --- a/cpp/src/io/parquet/stats_filter_helpers.hpp +++ b/cpp/src/io/parquet/stats_filter_helpers.hpp @@ -287,8 +287,10 @@ class stats_caster_base { auto [d_chars, d_offsets, _] = make_strings_children(val, chars, stream, mr); auto null_mask_buffer = cudf::create_null_mask(val.size(), cudf::mask_state::UNINITIALIZED, stream, mr); - CUDF_CUDA_TRY(cudf::detail::memcpy_async( - null_mask_buffer.data(), null_mask.data(), null_mask_buffer.size(), stream)); + CUDF_CUDA_TRY(cudf::detail::memcpy_async(null_mask_buffer.data(), + null_mask.data(), + null_mask_buffer.size() * sizeof(bitmask_type), + stream)); stream.sync(); return cudf::make_strings_column( val.size(), @@ -303,8 +305,10 @@ class stats_caster_base { auto data = cudf::detail::make_device_uvector_async(val, stream, mr); auto null_mask_buffer = cudf::create_null_mask(val.size(), cudf::mask_state::UNINITIALIZED, stream, mr); - CUDF_CUDA_TRY(cudf::detail::memcpy_async( - null_mask_buffer.data(), null_mask.data(), null_mask_buffer.size(), stream)); + CUDF_CUDA_TRY(cudf::detail::memcpy_async(null_mask_buffer.data(), + null_mask.data(), + null_mask_buffer.size() * sizeof(bitmask_type), + stream)); stream.sync(); return std::make_unique( dtype, val.size(), data.release(), std::move(null_mask_buffer), null_count); diff --git a/cpp/src/io/utilities/column_buffer.hpp b/cpp/src/io/utilities/column_buffer.hpp index caa3b3d5c6ff..6fe53dcb6bdf 100644 --- a/cpp/src/io/utilities/column_buffer.hpp +++ b/cpp/src/io/utilities/column_buffer.hpp @@ -128,14 +128,13 @@ class column_buffer_base { // input column, including same type, nullability, name, and user_data. static string_policy empty_like(string_policy const& input); - void set_null_mask(cuda::device_buffer&& mask) { _null_mask = std::move(mask); } - - template - auto null_mask() + void set_null_mask(cuda::device_buffer&& mask) { - return reinterpret_cast(_null_mask.data()); + _null_mask = std::move(mask); } - auto null_mask_size() { return _null_mask.size(); } + + auto null_mask() { return _null_mask.data(); } + auto null_mask_size() { return _null_mask.size() * sizeof(bitmask_type); } auto& null_count() { return _null_count; } auto data() { return static_cast(this)->data_impl(); } @@ -152,7 +151,7 @@ class column_buffer_base { protected: rmm::device_buffer _data{}; - cuda::device_buffer _null_mask = + cuda::device_buffer _null_mask = cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); size_type _null_count{0}; rmm::device_async_resource_ref _mr{cudf::get_current_device_resource_ref()}; diff --git a/cpp/src/io/utilities/data_casting.cu b/cpp/src/io/utilities/data_casting.cu index 07c359a5e21d..41d03a85336a 100644 --- a/cpp/src/io/utilities/data_casting.cu +++ b/cpp/src/io/utilities/data_casting.cu @@ -793,7 +793,7 @@ struct to_string_view_pair { template static std::unique_ptr parse_string(string_view_pair_it str_tuples, size_type col_size, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, cudf::detail::device_scalar& d_null_count, cudf::io::parse_options_view const& options, cuda::stream_ref stream, @@ -813,12 +813,8 @@ static std::unique_ptr parse_string(string_view_pair_it str_tuples, auto d_sizes = sizes.data(); auto null_count_data = d_null_count.data(); - auto single_thread_fn = - string_parse{str_tuples, - reinterpret_cast(null_mask.data()), - null_count_data, - options, - d_sizes}; + auto single_thread_fn = string_parse{ + str_tuples, null_mask.data(), null_count_data, options, d_sizes}; thrust::for_each_n(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), cuda::counting_iterator{0}, col_size, @@ -833,16 +829,15 @@ static std::unique_ptr parse_string(string_view_pair_it str_tuples, // TODO run these independent kernels in parallel streams. if (max_length > SINGLE_THREAD_THRESHOLD) { parse_fn_string_parallel - <<>>( - str_tuples, - col_size, - str_counter.data(), - reinterpret_cast(null_mask.data()), - null_count_data, - options, - d_sizes, - cudf::detail::input_offsetalator{}, - nullptr); + <<>>(str_tuples, + col_size, + str_counter.data(), + null_mask.data(), + null_count_data, + options, + d_sizes, + cudf::detail::input_offsetalator{}, + nullptr); CUDF_CUDA_TRY(cudaGetLastError()); } @@ -850,16 +845,15 @@ static std::unique_ptr parse_string(string_view_pair_it str_tuples, // for strings longer than WARP_THRESHOLD, 1 block per string str_counter.set_value(0, stream); parse_fn_string_parallel - <<>>( - str_tuples, - col_size, - str_counter.data(), - reinterpret_cast(null_mask.data()), - null_count_data, - options, - d_sizes, - cudf::detail::input_offsetalator{}, - nullptr); + <<>>(str_tuples, + col_size, + str_counter.data(), + null_mask.data(), + null_count_data, + options, + d_sizes, + cudf::detail::input_offsetalator{}, + nullptr); CUDF_CUDA_TRY(cudaGetLastError()); } @@ -881,16 +875,15 @@ static std::unique_ptr parse_string(string_view_pair_it str_tuples, if (max_length > SINGLE_THREAD_THRESHOLD) { str_counter.set_value(0, stream); parse_fn_string_parallel - <<>>( - str_tuples, - col_size, - str_counter.data(), - reinterpret_cast(null_mask.data()), - null_count_data, - options, - d_sizes, - d_offsets, - d_chars); + <<>>(str_tuples, + col_size, + str_counter.data(), + null_mask.data(), + null_count_data, + options, + d_sizes, + d_offsets, + d_chars); CUDF_CUDA_TRY(cudaGetLastError()); } @@ -898,16 +891,15 @@ static std::unique_ptr parse_string(string_view_pair_it str_tuples, str_counter.set_value(0, stream); // for strings longer than WARP_THRESHOLD, 1 block per string parse_fn_string_parallel - <<>>( - str_tuples, - col_size, - str_counter.data(), - reinterpret_cast(null_mask.data()), - null_count_data, - options, - d_sizes, - d_offsets, - d_chars); + <<>>(str_tuples, + col_size, + str_counter.data(), + null_mask.data(), + null_count_data, + options, + d_sizes, + d_offsets, + d_chars); CUDF_CUDA_TRY(cudaGetLastError()); } @@ -923,7 +915,7 @@ std::unique_ptr parse_data( cuda::zip_iterator offset_length_begin, size_type col_size, data_type col_type, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, size_type null_count, cudf::io::parse_options_view const& options, cuda::stream_ref stream, diff --git a/cpp/src/io/utilities/string_parsing.hpp b/cpp/src/io/utilities/string_parsing.hpp index 542cf313fad5..4cbfa07bfc0b 100644 --- a/cpp/src/io/utilities/string_parsing.hpp +++ b/cpp/src/io/utilities/string_parsing.hpp @@ -144,7 +144,7 @@ CUDF_EXPORT std::unique_ptr parse_data( cuda::zip_iterator offset_length_begin, size_type col_size, data_type col_type, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, size_type null_count, cudf::io::parse_options_view const& options, cuda::stream_ref stream, diff --git a/cpp/src/join/distinct_hash_join.cu b/cpp/src/join/distinct_hash_join.cu index a5a69a9b786f..07a37fcc3974 100644 --- a/cpp/src/join/distinct_hash_join.cu +++ b/cpp/src/join/distinct_hash_join.cu @@ -131,8 +131,7 @@ void find_matches_in_hash_table(HashTableType const& hash_table, auto stencil = cuda::counting_iterator{0}; auto const row_bitmask = cudf::detail::bitmask_and(left, stream, cudf::get_current_device_resource_ref()).first; - auto const pred = - cudf::detail::row_is_valid{reinterpret_cast(row_bitmask.data())}; + auto const pred = cudf::detail::row_is_valid{row_bitmask.data()}; hash_table.find_if_async( iter, iter + left_table_num_rows, stencil, pred, d_equal, hasher, found_begin, stream.get()); @@ -183,8 +182,7 @@ distinct_hash_join::distinct_hash_join(cudf::table_view const& right, auto stencil = cuda::counting_iterator{0}; auto const row_bitmask = cudf::detail::bitmask_and(_right, stream, cudf::get_current_device_resource_ref()).first; - auto const pred = - cudf::detail::row_is_valid{reinterpret_cast(row_bitmask.data())}; + auto const pred = cudf::detail::row_is_valid{row_bitmask.data()}; // insert valid rows this->_hash_table.insert_if_async( diff --git a/cpp/src/join/filtered_join/filtered_join.cu b/cpp/src/join/filtered_join/filtered_join.cu index 8a3e9150f5e4..c68efb854e19 100644 --- a/cpp/src/join/filtered_join/filtered_join.cu +++ b/cpp/src/join/filtered_join/filtered_join.cu @@ -41,8 +41,10 @@ namespace detail { /** * @brief Returns a validity mask for rows without nulls at any nested level, or null when unused. */ -std::pair, bitmask_type const*> make_filtered_join_row_bitmask( - table_view const& input, null_equality nulls_equal, cuda::stream_ref stream) +std::pair, bitmask_type const*> +make_filtered_join_row_bitmask(table_view const& input, + null_equality nulls_equal, + cuda::stream_ref stream) { if (nulls_equal == null_equality::EQUAL || !has_nested_nulls(input)) { return std::pair(cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream), nullptr); @@ -56,7 +58,7 @@ std::pair, bitmask_type const*> make_filtered_joi cudf::detail::bitmask_and( table_view{nullable_columns}, stream, cudf::get_current_device_resource_ref()) .first; - auto const row_bitmask_ptr = reinterpret_cast(row_bitmask.data()); + auto const row_bitmask_ptr = row_bitmask.data(); return std::pair(std::move(row_bitmask), row_bitmask_ptr); } diff --git a/cpp/src/join/filtered_join/filtered_join_common.cuh b/cpp/src/join/filtered_join/filtered_join_common.cuh index 166d7cb6566e..243a05d9e9d3 100644 --- a/cpp/src/join/filtered_join/filtered_join_common.cuh +++ b/cpp/src/join/filtered_join/filtered_join_common.cuh @@ -36,8 +36,10 @@ struct precomputed_hash { hash_value_type const* _hashes; }; -std::pair, bitmask_type const*> make_filtered_join_row_bitmask( - table_view const& input, null_equality nulls_equal, cuda::stream_ref stream); +std::pair, bitmask_type const*> +make_filtered_join_row_bitmask(table_view const& input, + null_equality nulls_equal, + cuda::stream_ref stream); class filtered_join_row_is_valid { public: diff --git a/cpp/src/join/hash_join/hash_join.cu b/cpp/src/join/hash_join/hash_join.cu index fed413f89867..eb4cbcb1b9d4 100644 --- a/cpp/src/join/hash_join/hash_join.cu +++ b/cpp/src/join/hash_join/hash_join.cu @@ -110,9 +110,7 @@ hash_join::hash_join(cudf::table_view const& right, _nulls_equal == null_equality::UNEQUAL ? cudf::detail::bitmask_and(right, stream, temp_mr).first : cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream, temp_mr); - auto const valid_rows = _nulls_equal == null_equality::UNEQUAL - ? reinterpret_cast(row_bitmask.data()) - : nullptr; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL ? row_bitmask.data() : nullptr; // Hashing and comparing variable-width rows again can dominate construction. Cache one // representative index per row for keys containing lists or strings, including within structs. auto const cache_representatives = std::any_of(right.begin(), right.end(), has_list_or_string); diff --git a/cpp/src/join/hash_join/hash_join_retrieve.cu b/cpp/src/join/hash_join/hash_join_retrieve.cu index 96a0d0aa666b..8f48d01d248c 100644 --- a/cpp/src/join/hash_join/hash_join_retrieve.cu +++ b/cpp/src/join/hash_join/hash_join_retrieve.cu @@ -92,9 +92,7 @@ hash_join::join_retrieve(cudf::table_view const& left, auto matched_build_rows = cudf::detail::device_scalar(0, stream, temp_mr); auto const row_bitmask = cudf::detail::bitmask_and(left, stream, temp_mr).first; - auto const valid_rows = _nulls_equal == null_equality::UNEQUAL - ? reinterpret_cast(row_bitmask.data()) - : nullptr; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL ? row_bitmask.data() : nullptr; auto count_matches = [&](auto equality, auto hasher) { launch_hash_csr_probe_count_kernel( diff --git a/cpp/src/join/hash_join/hash_join_size.cu b/cpp/src/join/hash_join/hash_join_size.cu index e292b612cb5e..e8bbb03286eb 100644 --- a/cpp/src/join/hash_join/hash_join_size.cu +++ b/cpp/src/join/hash_join/hash_join_size.cu @@ -47,9 +47,7 @@ std::size_t hash_join::join_size(cudf::table_view const& left, auto match_counts = cudf::detail::make_zeroed_device_buffer_async(left.num_rows(), stream, temp_mr); auto const row_bitmask = cudf::detail::bitmask_and(left, stream, temp_mr).first; - auto const valid_rows = _nulls_equal == null_equality::UNEQUAL - ? reinterpret_cast(row_bitmask.data()) - : nullptr; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL ? row_bitmask.data() : nullptr; auto count_matches = [&](auto equality, auto hasher) { launch_hash_csr_probe_count_kernel(left.num_rows(), @@ -96,9 +94,7 @@ std::size_t hash_join::join_size(cudf::table_view const& left, _right.num_rows(), stream, temp_mr); auto matched_build_rows = cudf::detail::device_scalar(0, stream, temp_mr); auto const row_bitmask = cudf::detail::bitmask_and(left, stream, temp_mr).first; - auto const valid_rows = _nulls_equal == null_equality::UNEQUAL - ? reinterpret_cast(row_bitmask.data()) - : nullptr; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL ? row_bitmask.data() : nullptr; auto count_matches = [&](auto equality, auto hasher) { launch_hash_csr_probe_count_kernel(left.num_rows(), diff --git a/cpp/src/join/hash_join/match_context.cu b/cpp/src/join/hash_join/match_context.cu index 570b8afc123b..e6b7b37df77c 100644 --- a/cpp/src/join/hash_join/match_context.cu +++ b/cpp/src/join/hash_join/match_context.cu @@ -42,9 +42,7 @@ std::unique_ptr> hash_join::make_match_co left, stream, cudf::get_current_device_resource_ref()); auto const temp_mr = cudf::get_current_device_resource_ref(); auto const row_bitmask = cudf::detail::bitmask_and(left, stream, temp_mr).first; - auto const valid_rows = _nulls_equal == null_equality::UNEQUAL - ? reinterpret_cast(row_bitmask.data()) - : nullptr; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL ? row_bitmask.data() : nullptr; auto count_matches = [&](auto equality, auto hasher) { if (join == join_kind::INNER_JOIN) { diff --git a/cpp/src/join/hash_join/partitioned_join_retrieve.cu b/cpp/src/join/hash_join/partitioned_join_retrieve.cu index c68a83a280e8..db20f824d032 100644 --- a/cpp/src/join/hash_join/partitioned_join_retrieve.cu +++ b/cpp/src/join/hash_join/partitioned_join_retrieve.cu @@ -110,9 +110,7 @@ hash_join::partitioned_join_retrieve(join_kind join, cuda::device_buffer probe_groups(stream, temp_mr, partition_size, cuda::no_init); auto const row_bitmask = cudf::detail::bitmask_and(left_partition_view, stream, temp_mr).first; - auto const valid_rows = _nulls_equal == null_equality::UNEQUAL - ? reinterpret_cast(row_bitmask.data()) - : nullptr; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL ? row_bitmask.data() : nullptr; auto save_groups = [&](auto equality, auto hasher) { if (join == join_kind::INNER_JOIN) { launch_hash_csr_probe_count_kernel(partition_size, diff --git a/cpp/src/join/key_remapping.cu b/cpp/src/join/key_remapping.cu index cfa283f01247..824490c4387e 100644 --- a/cpp/src/join/key_remapping.cu +++ b/cpp/src/join/key_remapping.cu @@ -318,8 +318,7 @@ class key_remap_table : public key_remap_table_interface { skip_nulls ? cudf::detail::bitmask_and(_right, stream, cudf::get_current_device_resource_ref()).first : cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); - auto const bitmask_ptr = - skip_nulls ? reinterpret_cast(row_bitmask.data()) : nullptr; + auto const bitmask_ptr = skip_nulls ? row_bitmask.data() : nullptr; if (compute_metrics) { // Use hash-based atomic counting for metrics computation @@ -468,8 +467,7 @@ class key_remap_table : public key_remap_table_interface { auto stencil = cuda::counting_iterator{0}; auto const row_bitmask = cudf::detail::bitmask_and(left_keys, stream, cudf::get_current_device_resource_ref()).first; - auto const pred = - row_is_valid{reinterpret_cast(row_bitmask.data())}; + auto const pred = row_is_valid{row_bitmask.data()}; _hash_table.find_if_async(iter, iter + left_num_rows, diff --git a/cpp/src/join/mark_join.cu b/cpp/src/join/mark_join.cu index 47d288564c87..b1e065887957 100644 --- a/cpp/src/join/mark_join.cu +++ b/cpp/src/join/mark_join.cu @@ -40,7 +40,7 @@ namespace cudf { namespace detail { namespace { -std::pair, bitmask_type const*> build_row_bitmask( +std::pair, bitmask_type const*> build_row_bitmask( table_view const& input, cuda::stream_ref stream) { auto const nullable_columns = get_nullable_columns(input); @@ -51,7 +51,7 @@ std::pair, bitmask_type const*> build_row_bitmask if (nullable_columns.size() > 1) { auto row_bitmask = cudf::detail::bitmask_and(table_view{nullable_columns}, stream, temp_mr).first; - auto const row_bitmask_ptr = reinterpret_cast(row_bitmask.data()); + auto const row_bitmask_ptr = row_bitmask.data(); return std::pair(std::move(row_bitmask), row_bitmask_ptr); } diff --git a/cpp/src/join/mixed_join_semi.cu b/cpp/src/join/mixed_join_semi.cu index 45b3c337ffdb..1e87c4025732 100644 --- a/cpp/src/join/mixed_join_semi.cu +++ b/cpp/src/join/mixed_join_semi.cu @@ -164,7 +164,7 @@ std::unique_ptr> mixed_join_semi( } else { cuda::counting_iterator stencil(0); auto const [row_bitmask, _] = cudf::detail::bitmask_and(right, stream, temp_mr); - row_is_valid pred{reinterpret_cast(row_bitmask.data())}; + row_is_valid pred{row_bitmask.data()}; // insert valid rows row_set.insert_if_async(iter, iter + right_num_rows, stencil, pred, stream.get()); diff --git a/cpp/src/join/sort_merge_join.cu b/cpp/src/join/sort_merge_join.cu index a0ddf0915cf6..74cfa4c6d897 100644 --- a/cpp/src/join/sort_merge_join.cu +++ b/cpp/src/join/sort_merge_join.cu @@ -581,14 +581,13 @@ void sort_merge_join::preprocessed_table::populate_nonnull_filter(cuda::stream_r stream, temp_mr); - thrust::for_each( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::counting_iterator{0}, - cuda::counting_iterator{0} + subset_size, - list_nonnull_filter{reinterpret_cast(validity_mask.data()), - reinterpret_cast(reduced_validity_mask.data()), - child_positions, - static_cast(subset_offset)}); + thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator{0}, + cuda::counting_iterator{0} + subset_size, + list_nonnull_filter{validity_mask.data(), + reduced_validity_mask.data(), + child_positions, + static_cast(subset_offset)}); } else if (col.type().id() == type_id::STRUCT) { // Recursive lambda to traverse struct hierarchy and accumulate null information. // This lambda ANDs the column's null mask with the accumulated mask in-place and recursively @@ -600,8 +599,7 @@ void sort_merge_join::preprocessed_table::populate_nonnull_filter(cuda::stream_r if (colview.nullable()) { // AND this column's null mask with the accumulated mask auto colmask = colview.null_mask(); - std::vector masks{reinterpret_cast(colmask), - reinterpret_cast(mask)}; + std::vector masks{colmask, mask}; std::vector begin_bits{0, 0}; cudf::detail::inplace_bitmask_and( device_span(mask, num_bitmask_words(num_rows)), @@ -629,12 +627,11 @@ void sort_merge_join::preprocessed_table::populate_nonnull_filter(cuda::stream_r // Process all children of the struct column for (auto it = col.child_begin(); it != col.child_end(); it++) { auto& child = *it; - and_bitmasks(and_bitmasks, reinterpret_cast(validity_mask.data()), child); + and_bitmasks(and_bitmasks, validity_mask.data(), child); } } } - this->_num_nulls = - null_count(reinterpret_cast(validity_mask.data()), 0, table.num_rows(), stream); + this->_num_nulls = null_count(validity_mask.data(), 0, table.num_rows(), stream); this->_validity_mask = std::move(validity_mask); } @@ -727,13 +724,12 @@ cuda::device_buffer sort_merge_join::preprocessed_table::map_table_to auto temp_mr = cudf::get_current_device_resource_ref(); auto const table_mapping_size = _table_view.num_rows() - _num_nulls.value(); cuda::device_buffer table_mapping(stream, temp_mr, table_mapping_size, cuda::no_init); - cudf::detail::copy_if_async( - cuda::counting_iterator{0}, - cuda::counting_iterator{_table_view.num_rows()}, - cuda::counting_iterator{0}, - table_mapping.begin(), - is_row_valid{reinterpret_cast(_validity_mask.value().data())}, - stream); + cudf::detail::copy_if_async(cuda::counting_iterator{0}, + cuda::counting_iterator{_table_view.num_rows()}, + cuda::counting_iterator{0}, + table_mapping.begin(), + is_row_valid{_validity_mask.value().data()}, + stream); return table_mapping; } @@ -869,8 +865,7 @@ sort_merge_join::left_join(table_view const& left, auto const total_output_size = preprocessed_left_indices->size() + static_cast(num_filtered_nulls); - auto const validity_mask = - reinterpret_cast(preprocessed_left._validity_mask.value().data()); + auto const validity_mask = preprocessed_left._validity_mask.value().data(); cuda::device_buffer null_left_indices(stream, cudf::get_current_device_resource_ref(), static_cast(num_filtered_nulls), diff --git a/cpp/src/join/sort_merge_join.hpp b/cpp/src/join/sort_merge_join.hpp index ec7ade0b5593..2a112a7d9e99 100644 --- a/cpp/src/join/sort_merge_join.hpp +++ b/cpp/src/join/sort_merge_join.hpp @@ -120,7 +120,7 @@ class sort_merge_join { ///< rows of the unprocessed table view if null equality is set ///< to false, otherwise equal to the unprocessed table view - std::optional> _validity_mask = + std::optional> _validity_mask = std::nullopt; ///< Optional validity mask for null_equality::UNEQUAL case std::optional _num_nulls = std::nullopt; ///< Optional count of nulls for null_equality::UNEQUAL case diff --git a/cpp/src/join/streaming_hash_join.cu b/cpp/src/join/streaming_hash_join.cu index fd081a08b5a1..8c6e4222dd9a 100644 --- a/cpp/src/join/streaming_hash_join.cu +++ b/cpp/src/join/streaming_hash_join.cu @@ -358,7 +358,7 @@ class streaming_hash_join_impl { auto preprocessed = row::equality::preprocessed_table::create(keys, stream, mr.get_output_mr()); auto const batch_rows = keys.num_rows(); - auto row_bitmask = [&]() -> std::optional> { + auto row_bitmask = [&]() -> std::optional> { if (batch_rows > 0 && compare_nulls == null_equality::UNEQUAL && nullable(keys)) { return cudf::detail::bitmask_and(keys, stream, mr.get_temporary_mr()).first; } @@ -395,12 +395,11 @@ class streaming_hash_join_impl { if (compare_nulls == null_equality::EQUAL || !nullable(keys)) { hash_table.insert_async(input_begin, input_begin + batch_rows, stream.get()); } else { - hash_table.insert_if_async( - input_begin, - input_begin + batch_rows, - cuda::counting_iterator{0}, - row_is_valid{reinterpret_cast(row_bitmask->data())}, - stream.get()); + hash_table.insert_if_async(input_begin, + input_begin + batch_rows, + cuda::counting_iterator{0}, + row_is_valid{row_bitmask->data()}, + stream.get()); } }; diff --git a/cpp/src/json/json_path.cu b/cpp/src/json/json_path.cu index d43fccedf5d2..7a202a5d0e5d 100644 --- a/cpp/src/json/json_path.cu +++ b/cpp/src/json/json_path.cu @@ -1022,7 +1022,7 @@ std::unique_ptr get_json_object(cudf::strings_column_view const& c // potential optimization : if we know that all outputs are valid, we could skip creating // the validity mask altogether - cuda::device_buffer validity = + cuda::device_buffer validity = cudf::detail::create_null_mask(col.size(), mask_state::UNINITIALIZED, stream, mr); // compute results @@ -1036,7 +1036,7 @@ std::unique_ptr get_json_object(cudf::strings_column_view const& c sizes.data(), d_offsets, chars.data(), - reinterpret_cast(validity.data()), + validity.data(), d_valid_count.data(), options); CUDF_CUDA_TRY(cudaGetLastError()); diff --git a/cpp/src/lists/combine/concatenate_rows.cu b/cpp/src/lists/combine/concatenate_rows.cu index a0f45a0058d1..5fff0ce91507 100644 --- a/cpp/src/lists/combine/concatenate_rows.cu +++ b/cpp/src/lists/combine/concatenate_rows.cu @@ -59,7 +59,7 @@ namespace { * We can do this by recomputing a new offsets column that does this regrouping. * */ -std::tuple, cuda::device_buffer, size_type> +std::tuple, cuda::device_buffer, size_type> generate_regrouped_offsets_and_null_mask(table_device_view const& input, bool build_null_mask, concatenate_null_policy null_policy, @@ -125,7 +125,7 @@ generate_regrouped_offsets_and_null_mask(table_device_view const& input, auto [null_mask, null_count] = [&]() { // if the input doesn't contain nulls, no work to do if (!build_null_mask) { - return std::pair, size_type>{ + return std::pair, size_type>{ cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), 0}; } diff --git a/cpp/src/lists/copying/concatenate.cu b/cpp/src/lists/copying/concatenate.cu index 9947fa3c85cf..48d2e0d36221 100644 --- a/cpp/src/lists/copying/concatenate.cu +++ b/cpp/src/lists/copying/concatenate.cu @@ -115,9 +115,9 @@ std::unique_ptr concatenate(host_span columns, // if any of the input columns have nulls, construct the output mask bool const has_nulls = std::any_of(columns.begin(), columns.end(), [](auto const& col) { return col.has_nulls(); }); - cuda::device_buffer null_mask = cudf::detail::create_null_mask( + cuda::device_buffer null_mask = cudf::detail::create_null_mask( total_list_count, has_nulls ? mask_state::UNINITIALIZED : mask_state::UNALLOCATED, stream, mr); - auto null_mask_data = reinterpret_cast(null_mask.data()); + auto null_mask_data = null_mask.data(); auto const null_count = has_nulls ? cudf::detail::concatenate_masks(columns, null_mask_data, stream) : size_type{0}; diff --git a/cpp/src/lists/copying/copying.cu b/cpp/src/lists/copying/copying.cu index 0e5e66a022d0..f75344d6d301 100644 --- a/cpp/src/lists/copying/copying.cu +++ b/cpp/src/lists/copying/copying.cu @@ -75,8 +75,7 @@ std::unique_ptr copy_slice(lists_column_view const& lists, // Compute the null mask of the result: auto null_mask = cudf::detail::copy_bitmask(lists.null_mask(), start, end, stream, mr); - auto null_count = cudf::detail::null_count( - reinterpret_cast(null_mask.data()), 0, end - start, stream); + auto null_count = cudf::detail::null_count(null_mask.data(), 0, end - start, stream); return make_lists_column( lists_count, std::move(offsets), std::move(child), null_count, std::move(null_mask)); diff --git a/cpp/src/lists/copying/scatter_helper.cu b/cpp/src/lists/copying/scatter_helper.cu index c6268d5a641c..935a374ef2cb 100644 --- a/cpp/src/lists/copying/scatter_helper.cu +++ b/cpp/src/lists/copying/scatter_helper.cu @@ -35,10 +35,10 @@ namespace detail { * @param num_child_rows Number of rows in child column * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate child column's null mask - * @return std::pair, size_type> Child column's null mask and null - * row count + * @return std::pair, size_type> Child column's null mask + * and null row count */ -std::pair, size_type> construct_child_nullmask( +std::pair, size_type> construct_child_nullmask( rmm::device_uvector const& parent_list_vector, column_view const& parent_list_offsets, cudf::lists_column_device_view const& source_lists, diff --git a/cpp/src/lists/explode.cu b/cpp/src/lists/explode.cu index 46eebdef657d..594eb0b97b7b 100644 --- a/cpp/src/lists/explode.cu +++ b/cpp/src/lists/explode.cu @@ -88,7 +88,7 @@ std::unique_ptr
build_table( [] __device__(auto i) { return i != InvalidIndex; }, stream, mr) - : std::pair, size_type>{ + : std::pair, size_type>{ cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED, stream), size_type{0}}; columns.insert(columns.begin() + explode_column_idx, diff --git a/cpp/src/lists/lists_column_factories.cu b/cpp/src/lists/lists_column_factories.cu index f1035ec43948..fd38f93607db 100644 --- a/cpp/src/lists/lists_column_factories.cu +++ b/cpp/src/lists/lists_column_factories.cu @@ -53,13 +53,8 @@ std::unique_ptr make_lists_column_from_scalar(list_scalar const& v } auto children_views = std::vector{offsets->view(), value.view()}; - auto one_row_col_view = column_view(data_type{type_id::LIST}, - 1, - nullptr, - reinterpret_cast(null_mask.data()), - null_count, - 0, - children_views); + auto one_row_col_view = column_view( + data_type{type_id::LIST}, 1, nullptr, null_mask.data(), null_count, 0, children_views); auto begin = cuda::make_constant_iterator(0); auto res = cudf::detail::gather(table_view({one_row_col_view}), @@ -113,7 +108,7 @@ std::unique_ptr make_lists_column(size_type num_rows, std::unique_ptr offsets_column, std::unique_ptr child_column, size_type null_count, - cuda::device_buffer&& null_mask) + cuda::device_buffer&& null_mask) { if (null_count > 0) { CUDF_EXPECTS(null_mask.size() > 0, "Column with nulls must be nullable."); } CUDF_EXPECTS( diff --git a/cpp/src/merge/merge.cu b/cpp/src/merge/merge.cu index 7d86d7e3efa6..49688b3c9cb3 100644 --- a/cpp/src/merge/merge.cu +++ b/cpp/src/merge/merge.cu @@ -503,17 +503,12 @@ std::unique_ptr column_merger::operator()( auto const merged_size = lcol.size() + rcol.size(); // materialize the output buffer - cuda::device_buffer validity = + cuda::device_buffer validity = lcol.has_nulls() || rcol.has_nulls() ? detail::create_null_mask(merged_size, mask_state::UNINITIALIZED, stream, mr) : cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); if (lcol.has_nulls() || rcol.has_nulls()) { - materialize_bitmask(lcol, - rcol, - reinterpret_cast(validity.data()), - merged_size, - row_order_.data(), - stream); + materialize_bitmask(lcol, rcol, validity.data(), merged_size, row_order_.data(), stream); } return make_structs_column(merged_size, diff --git a/cpp/src/quantiles/tdigest/tdigest.cu b/cpp/src/quantiles/tdigest/tdigest.cu index 0cc2b6eb7a6a..3b62ea818d81 100644 --- a/cpp/src/quantiles/tdigest/tdigest.cu +++ b/cpp/src/quantiles/tdigest/tdigest.cu @@ -227,7 +227,7 @@ std::unique_ptr compute_approx_percentiles(tdigest_column_view const& in }, stream, mr) - : std::pair, size_type>{ + : std::pair, size_type>{ cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), 0}; }(); @@ -385,7 +385,7 @@ std::unique_ptr percentile_approx(tdigest_column_view const& input, tdigest_is_empty + tdv.size(), 0); if (null_count == 0) { - return std::pair, size_type>{ + return std::pair, size_type>{ cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), null_count}; } return cudf::detail::valid_if( diff --git a/cpp/src/reductions/approx_distinct_count.cu b/cpp/src/reductions/approx_distinct_count.cu index 51a8416bc085..aefbb135f98c 100644 --- a/cpp/src/reductions/approx_distinct_count.cu +++ b/cpp/src/reductions/approx_distinct_count.cu @@ -247,8 +247,7 @@ void approx_distinct_count::add(table_view const& input, cuda::stream_re ref.add_async(hash_iter, hash_iter + num_rows, stream); } else { auto const row_bitmask = cudf::detail::bitmask_and(input, stream, temp_mr).first; - auto const pred = - row_is_valid{reinterpret_cast(row_bitmask.data())}; + auto const pred = row_is_valid{row_bitmask.data()}; ref.add_if_async(hash_iter, hash_iter + num_rows, stencil, pred, stream); } } else { @@ -258,7 +257,7 @@ void approx_distinct_count::add(table_view const& input, cuda::stream_re ref.add_if_async(hash_iter, hash_iter + num_rows, stencil, pred, stream); } else { auto const row_bitmask = cudf::detail::bitmask_and(input, stream, temp_mr).first; - auto const bitmask_ptr = reinterpret_cast(row_bitmask.data()); + auto const bitmask_ptr = row_bitmask.data(); auto const pred = check_nans_predicate{*d_table, bitmask_ptr}; ref.add_if_async(hash_iter, hash_iter + num_rows, stencil, pred, stream); } diff --git a/cpp/src/reductions/distinct_count.cu b/cpp/src/reductions/distinct_count.cu index 2689b79be96f..331f0d1a553c 100644 --- a/cpp/src/reductions/distinct_count.cu +++ b/cpp/src/reductions/distinct_count.cu @@ -168,7 +168,7 @@ cudf::size_type distinct_count(table_view const& keys, // zero. So, it is only when null_count is not zero that we need // to do a filtered insertion. if (null_count > 0) { - row_validity pred{reinterpret_cast(row_bitmask.data())}; + row_validity pred{row_bitmask.data()}; return key_set.insert_if(iter, iter + num_rows, stencil, pred, stream.get()) + 1; } } diff --git a/cpp/src/reductions/scan/scan.cuh b/cpp/src/reductions/scan/scan.cuh index 4ca917008d25..2b2b655f6c2a 100644 --- a/cpp/src/reductions/scan/scan.cuh +++ b/cpp/src/reductions/scan/scan.cuh @@ -20,10 +20,11 @@ namespace cudf { namespace detail { // logical-and scan of the null mask of the input view -std::pair, size_type> mask_scan(column_view const& input_view, - scan_type inclusive, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr); +std::pair, size_type> mask_scan( + column_view const& input_view, + scan_type inclusive, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr); // exponentially weighted moving average of the input std::unique_ptr exponentially_weighted_moving_average(column_view const& input, diff --git a/cpp/src/reductions/scan/scan_exclusive.cu b/cpp/src/reductions/scan/scan_exclusive.cu index 5b7ba872f822..c53f390005a6 100644 --- a/cpp/src/reductions/scan/scan_exclusive.cu +++ b/cpp/src/reductions/scan/scan_exclusive.cu @@ -96,8 +96,7 @@ std::unique_ptr scan_exclusive(column_view const& input, return std::make_pair(cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), size_type{0}); }(); - auto output = scan_agg_dispatch( - input, agg, reinterpret_cast(mask.data()), stream, mr); + auto output = scan_agg_dispatch(input, agg, mask.data(), stream, mr); output->set_null_mask(std::move(mask), null_count); return output; diff --git a/cpp/src/reductions/scan/scan_inclusive.cu b/cpp/src/reductions/scan/scan_inclusive.cu index 2dd357f2cca5..fe674ab2382b 100644 --- a/cpp/src/reductions/scan/scan_inclusive.cu +++ b/cpp/src/reductions/scan/scan_inclusive.cu @@ -33,12 +33,13 @@ namespace cudf { namespace detail { // logical-and scan of the null mask of the input view -std::pair, size_type> mask_scan(column_view const& input_view, - scan_type inclusive, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) +std::pair, size_type> mask_scan( + column_view const& input_view, + scan_type inclusive, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) { - cuda::device_buffer mask = + cuda::device_buffer mask = detail::create_null_mask(input_view.size(), mask_state::UNINITIALIZED, stream, mr); auto d_input = column_device_view::create(input_view, stream); auto valid_itr = detail::make_validity_iterator(*d_input); @@ -54,12 +55,8 @@ std::pair, size_type> mask_scan(column_view const return std::min(input_view.size(), first_null + exclusive_offset); }(); - set_null_mask(reinterpret_cast(mask.data()), 0, first_null_position, true, stream); - set_null_mask(reinterpret_cast(mask.data()), - first_null_position, - input_view.size(), - false, - stream); + set_null_mask(mask.data(), 0, first_null_position, true, stream); + set_null_mask(mask.data(), first_null_position, input_view.size(), false, stream); return {std::move(mask), input_view.size() - first_null_position}; } @@ -213,8 +210,7 @@ std::unique_ptr scan_inclusive(column_view const& input, return std::make_pair(cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), size_type{0}); }(); - auto output = scan_agg_dispatch( - input, agg, reinterpret_cast(mask.data()), stream, mr); + auto output = scan_agg_dispatch(input, agg, mask.data(), stream, mr); // Use the null mask produced by the op for EWM if (agg.kind != aggregation::EWMA) { output->set_null_mask(std::move(mask), null_count); } @@ -226,7 +222,7 @@ std::unique_ptr scan_inclusive(column_view const& input, auto content = output->release(); // Build new children columns. - auto const null_mask = reinterpret_cast(content.null_mask->data()); + auto const null_mask = content.null_mask->data(); std::for_each(content.children.begin(), content.children.end(), [null_mask, null_count, stream, mr](auto& child) { diff --git a/cpp/src/scalar/scalar.cpp b/cpp/src/scalar/scalar.cpp index 52c2b5820e0c..bf043dae0fc5 100644 --- a/cpp/src/scalar/scalar.cpp +++ b/cpp/src/scalar/scalar.cpp @@ -628,7 +628,7 @@ table struct_scalar::init_data(table&& data, 1, mask_state::ALL_NULL, stream, cudf::get_current_device_resource_ref()); for (auto& col : data_cols) { col = cudf::structs::detail::superimpose_and_sanitize_nulls( - reinterpret_cast(validity.data()), 1, std::move(col), stream, mr); + validity.data(), 1, std::move(col), stream, mr); } return table{std::move(data_cols)}; diff --git a/cpp/src/search/contains_table_impl.cu b/cpp/src/search/contains_table_impl.cu index 1f5ab9ed5280..fd0b65fe11a8 100644 --- a/cpp/src/search/contains_table_impl.cu +++ b/cpp/src/search/contains_table_impl.cu @@ -19,7 +19,7 @@ namespace cudf::detail { * @param stream CUDA stream used for device memory operations and kernel launches * @return A pair of pointer to the output bitmask and the buffer containing the bitmask */ -std::pair, bitmask_type const*> build_row_bitmask( +std::pair, bitmask_type const*> build_row_bitmask( table_view const& input, cuda::stream_ref stream) { auto const nullable_columns = get_nullable_columns(input); @@ -33,7 +33,7 @@ std::pair, bitmask_type const*> build_row_bitmask cudf::detail::bitmask_and( table_view{nullable_columns}, stream, cudf::get_current_device_resource_ref()) .first; - auto const row_bitmask_ptr = reinterpret_cast(row_bitmask.data()); + auto const row_bitmask_ptr = row_bitmask.data(); return std::pair(std::move(row_bitmask), row_bitmask_ptr); } diff --git a/cpp/src/search/contains_table_impl.cuh b/cpp/src/search/contains_table_impl.cuh index 2c5788069581..e2a6849f4277 100644 --- a/cpp/src/search/contains_table_impl.cuh +++ b/cpp/src/search/contains_table_impl.cuh @@ -93,7 +93,7 @@ struct comparator_adapter { * @param stream CUDA stream used for device memory operations and kernel launches * @return A pair of pointer to the output bitmask and the buffer containing the bitmask */ -std::pair, bitmask_type const*> build_row_bitmask( +std::pair, bitmask_type const*> build_row_bitmask( table_view const& input, cuda::stream_ref stream); /** diff --git a/cpp/src/strings/char_types/char_types.cu b/cpp/src/strings/char_types/char_types.cu index 89a7a0ab95fb..6c27076cbf98 100644 --- a/cpp/src/strings/char_types/char_types.cu +++ b/cpp/src/strings/char_types/char_types.cu @@ -188,7 +188,7 @@ std::unique_ptr filter_characters_of_type(strings_column_view const& str d_replacement}; // copy null mask from input column - cuda::device_buffer null_mask = + cuda::device_buffer null_mask = cudf::detail::copy_bitmask(strings.parent(), stream, mr); // this utility calls filterer to build the offsets and chars columns diff --git a/cpp/src/strings/convert/convert_booleans.cu b/cpp/src/strings/convert/convert_booleans.cu index 45fc5b3b0484..0a9c390e6d45 100644 --- a/cpp/src/strings/convert/convert_booleans.cu +++ b/cpp/src/strings/convert/convert_booleans.cu @@ -127,7 +127,8 @@ std::unique_ptr from_booleans(column_view const& booleans, auto d_column = *column; // copy null mask - cuda::device_buffer null_mask = cudf::detail::copy_bitmask(booleans, stream, mr); + cuda::device_buffer null_mask = + cudf::detail::copy_bitmask(booleans, stream, mr); auto [offsets, chars] = make_strings_children(from_booleans_fn{d_column, d_true, d_false}, strings_count, stream, mr); diff --git a/cpp/src/strings/convert/convert_durations.cu b/cpp/src/strings/convert/convert_durations.cu index 38c72aad7b0c..7ff85b975d4a 100644 --- a/cpp/src/strings/convert/convert_durations.cu +++ b/cpp/src/strings/convert/convert_durations.cu @@ -403,7 +403,8 @@ struct dispatch_from_durations_fn { auto d_column = *column; // copy null mask - cuda::device_buffer null_mask = cudf::detail::copy_bitmask(durations, stream, mr); + cuda::device_buffer null_mask = + cudf::detail::copy_bitmask(durations, stream, mr); auto [offsets, chars] = make_strings_children(from_durations_fn{d_column, d_format_items, compiler.items_count()}, diff --git a/cpp/src/strings/convert/convert_floats.cu b/cpp/src/strings/convert/convert_floats.cu index 4d06a8f9456d..1126a1b93112 100644 --- a/cpp/src/strings/convert/convert_floats.cu +++ b/cpp/src/strings/convert/convert_floats.cu @@ -394,7 +394,8 @@ struct dispatch_from_floats_fn { auto d_column = *column; // copy the null mask - cuda::device_buffer null_mask = cudf::detail::copy_bitmask(floats, stream, mr); + cuda::device_buffer null_mask = + cudf::detail::copy_bitmask(floats, stream, mr); auto [offsets, chars] = make_strings_children(from_floats_fn{d_column}, strings_count, stream, mr); diff --git a/cpp/src/strings/convert/convert_integers.cu b/cpp/src/strings/convert/convert_integers.cu index 9e7bb608e2ed..a7df31a869e9 100644 --- a/cpp/src/strings/convert/convert_integers.cu +++ b/cpp/src/strings/convert/convert_integers.cu @@ -354,7 +354,8 @@ struct dispatch_from_integers_fn { auto d_column = *column; // copy the null mask - cuda::device_buffer null_mask = cudf::detail::copy_bitmask(integers, stream, mr); + cuda::device_buffer null_mask = + cudf::detail::copy_bitmask(integers, stream, mr); auto [offsets, chars] = make_strings_children(from_integers_fn{d_column}, strings_count, stream, mr); diff --git a/cpp/src/strings/convert/convert_urls.cu b/cpp/src/strings/convert/convert_urls.cu index 0a6bfc366cce..f62658558dd7 100644 --- a/cpp/src/strings/convert/convert_urls.cu +++ b/cpp/src/strings/convert/convert_urls.cu @@ -411,7 +411,7 @@ std::unique_ptr url_decode(strings_column_view const& strings, CUDF_CUDA_TRY(cudaGetLastError()); // copy null mask - cuda::device_buffer null_mask = + cuda::device_buffer null_mask = cudf::detail::copy_bitmask(strings.parent(), stream, mr); return make_strings_column(strings_count, diff --git a/cpp/src/strings/copying/concatenate.cu b/cpp/src/strings/copying/concatenate.cu index 09e339628e44..e186889146be 100644 --- a/cpp/src/strings/copying/concatenate.cu +++ b/cpp/src/strings/copying/concatenate.cu @@ -249,7 +249,7 @@ std::unique_ptr concatenate(host_span columns, static_cast(columns.size()), strings_count, itr_new_offsets, - reinterpret_cast(null_mask.data()), + null_mask.data(), d_valid_count.data()); CUDF_CUDA_TRY(cudaGetLastError()); diff --git a/cpp/src/strings/copying/copying.cu b/cpp/src/strings/copying/copying.cu index ea0cc16aa1c9..c1e6776aa012 100644 --- a/cpp/src/strings/copying/copying.cu +++ b/cpp/src/strings/copying/copying.cu @@ -67,8 +67,7 @@ std::unique_ptr copy_slice(strings_column_view const& input, auto null_mask = cudf::detail::copy_bitmask( input.null_mask(), offsets_offset, offsets_offset + strings_count, stream, mr); - auto null_count = cudf::detail::null_count( - reinterpret_cast(null_mask.data()), 0, strings_count, stream); + auto null_count = cudf::detail::null_count(null_mask.data(), 0, strings_count, stream); return make_strings_column(strings_count, std::move(offsets_column), diff --git a/cpp/src/strings/replace/replace.cu b/cpp/src/strings/replace/replace.cu index 6746eeae6664..f4fb00583b69 100644 --- a/cpp/src/strings/replace/replace.cu +++ b/cpp/src/strings/replace/replace.cu @@ -531,8 +531,7 @@ std::unique_ptr replace(strings_column_view const& input, return std::pair{cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), 0}; } }(); - auto const d_valid_mask = - has_any_nulls ? reinterpret_cast(null_mask.data()) : nullptr; + auto const d_valid_mask = has_any_nulls ? null_mask.data() : nullptr; auto d_targets = column_device_view::create(targets.parent(), stream); auto d_repls = column_device_view::create(repls.parent(), stream); diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu index 3b5cf1b5aa33..938d0423be7a 100644 --- a/cpp/src/strings/strings_column_factories.cu +++ b/cpp/src/strings/strings_column_factories.cu @@ -84,7 +84,7 @@ std::vector> make_strings_column_batch( auto [offsets_cols, d_chars_sizes] = make_offsets_child_column_batch_async(input, stream, mr); - std::vector> null_masks; + std::vector> null_masks; null_masks.reserve(num_columns); rmm::device_uvector d_valid_counts(num_columns, stream, mr); @@ -107,7 +107,7 @@ std::vector> make_strings_column_batch( cudf::detail::grid_1d{static_cast(string_count), block_size}; cudf::detail::valid_if_kernel <<>>( - reinterpret_cast(null_masks.back().data()), + null_masks.back().data(), string_pairs.data(), string_count, [] __device__(string_index_pair const pair) -> bool { return pair.first != nullptr; }, @@ -232,7 +232,7 @@ std::unique_ptr make_strings_column(size_type num_strings, std::unique_ptr offsets_column, rmm::device_buffer&& chars_buffer, size_type null_count, - cuda::device_buffer&& null_mask) + cuda::device_buffer&& null_mask) { CUDF_FUNC_RANGE(); @@ -255,7 +255,7 @@ std::unique_ptr make_strings_column(size_type num_strings, std::unique_ptr make_strings_column(size_type num_strings, rmm::device_uvector&& offsets, rmm::device_uvector&& chars, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, size_type null_count) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/strings/wrap.cu b/cpp/src/strings/wrap.cu index 83afa3159ed9..7580bb087dd7 100644 --- a/cpp/src/strings/wrap.cu +++ b/cpp/src/strings/wrap.cu @@ -97,7 +97,7 @@ std::unique_ptr wrap(strings_column_view const& strings, size_type null_count = strings.null_count(); // copy null mask - cuda::device_buffer null_mask = + cuda::device_buffer null_mask = cudf::detail::copy_bitmask(strings.parent(), stream, mr); // build offsets column diff --git a/cpp/src/structs/copying/concatenate.cu b/cpp/src/structs/copying/concatenate.cu index e2802ee13bda..3cd6c3d34cf2 100644 --- a/cpp/src/structs/copying/concatenate.cu +++ b/cpp/src/structs/copying/concatenate.cu @@ -53,12 +53,12 @@ std::unique_ptr concatenate(host_span columns, // if any of the input columns have nulls, construct the output mask bool const has_nulls = std::any_of(columns.begin(), columns.end(), [](auto const& col) { return col.has_nulls(); }); - cuda::device_buffer null_mask = + cuda::device_buffer null_mask = create_null_mask(total_length, has_nulls ? mask_state::UNINITIALIZED : mask_state::UNALLOCATED, stream, cudf::get_current_device_resource_ref()); - auto null_mask_data = reinterpret_cast(null_mask.data()); + auto null_mask_data = null_mask.data(); auto const null_count = has_nulls ? cudf::detail::concatenate_masks(columns, null_mask_data, stream) : size_type{0}; diff --git a/cpp/src/structs/structs_column_factories.cu b/cpp/src/structs/structs_column_factories.cu index 628baf3997f7..f2647878d81f 100644 --- a/cpp/src/structs/structs_column_factories.cu +++ b/cpp/src/structs/structs_column_factories.cu @@ -19,7 +19,7 @@ std::unique_ptr make_structs_column( size_type num_rows, std::vector>&& child_columns, size_type null_count, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { @@ -34,11 +34,7 @@ std::unique_ptr make_structs_column( if (null_mask.size() > 0) { for (auto& child : child_columns) { child = structs::detail::superimpose_and_sanitize_nulls( - reinterpret_cast(null_mask.data()), - null_count, - std::move(child), - stream, - mr); + null_mask.data(), null_count, std::move(child), stream, mr); } } @@ -55,14 +51,15 @@ std::unique_ptr create_structs_hierarchy( size_type num_rows, std::vector>&& child_columns, size_type null_count, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(null_count <= 0 || null_mask.size() > 0, "Struct column with nulls must be nullable."); - CUDF_EXPECTS(null_mask.size() == 0 || null_mask.size() == bitmask_allocation_size_bytes(num_rows), + CUDF_EXPECTS(null_mask.size() == 0 || + null_mask.size() * sizeof(bitmask_type) == bitmask_allocation_size_bytes(num_rows), "Number of bits in null_mask should equal number of rows in input columns"); CUDF_EXPECTS(std::all_of(child_columns.begin(), diff --git a/cpp/src/structs/utilities.cu b/cpp/src/structs/utilities.cu index e52376acf570..01d228eae4de 100644 --- a/cpp/src/structs/utilities.cu +++ b/cpp/src/structs/utilities.cu @@ -267,11 +267,8 @@ std::unique_ptr superimpose_nulls(bitmask_type const* null_mask, "Child columns must have the same number of rows as the Struct column."); for (auto& child : content.children) { - child = superimpose_nulls(reinterpret_cast(content.null_mask->data()), - new_null_count, - std::move(child), - stream, - mr); + child = + superimpose_nulls(content.null_mask->data(), new_null_count, std::move(child), stream, mr); } return std::make_unique(cudf::data_type{type_id::STRUCT}, num_rows, @@ -421,7 +418,7 @@ std::pair push_down_nulls_no_sanitize( auto parent_child_null_masks = std::vector{structs_view.null_mask(), child.null_mask()}; - auto [new_child_mask, null_count] = [&] { + auto [new_child_mask, null_count] = [&]() -> std::pair { if (not child.nullable()) { // Adopt parent STRUCT's null mask. return std::pair{structs_view.null_mask(), 0}; @@ -440,9 +437,7 @@ std::pair push_down_nulls_no_sanitize( stream, mr); ret_nullable_data.new_null_masks.push_back(std::move(new_mask)); - return std::pair{ - reinterpret_cast(ret_nullable_data.new_null_masks.back().data()), - null_count}; + return std::pair{ret_nullable_data.new_null_masks.back().data(), null_count}; }(); return column_view(child.type(), @@ -551,10 +546,11 @@ std::vector> enforce_null_consistency( // Helper struct to store and manipulate struct column properties struct contents { - size_type null_count = 0; // Number of null values - std::unique_ptr> null_mask = nullptr; // Null mask buffer - size_type num_children = 0; // Number of child columns - size_type num_elements = 0; // Number of rows + size_type null_count = 0; // Number of null values + std::unique_ptr> null_mask = + nullptr; // Null mask buffer + size_type num_children = 0; // Number of child columns + size_type num_elements = 0; // Number of rows }; std::vector struct_contents; // Store properties of each struct column @@ -573,10 +569,9 @@ std::vector> enforce_null_consistency( struct_contents.back().null_mask = std::move(col_contents.null_mask); // Add null mask pointers for each child of this struct - struct_root_masks.insert( - struct_root_masks.end(), - col_contents.children.size(), - reinterpret_cast(struct_contents.back().null_mask->data())); + struct_root_masks.insert(struct_root_masks.end(), + col_contents.children.size(), + struct_contents.back().null_mask->data()); // Add all child columns from this struct for (auto& child : col_contents.children) { diff --git a/cpp/src/transform/bools_to_mask.cu b/cpp/src/transform/bools_to_mask.cu index ee61586443e6..a0dbdf6a3fde 100644 --- a/cpp/src/transform/bools_to_mask.cu +++ b/cpp/src/transform/bools_to_mask.cu @@ -20,13 +20,13 @@ namespace cudf { namespace detail { -std::pair>, cudf::size_type> bools_to_mask( +std::pair>, cudf::size_type> bools_to_mask( column_view const& input, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.type().id() == type_id::BOOL8, "Input is not of type bool"); if (input.is_empty()) { - return std::pair(std::make_unique>( + return std::pair(std::make_unique>( cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)), 0); } @@ -40,20 +40,22 @@ std::pair>, cudf::size_type> bool auto mask = detail::valid_if(input_begin, input_begin + input.size(), pred, stream, mr); - return std::pair(std::make_unique>(std::move(mask.first)), - mask.second); + return std::pair( + std::make_unique>(std::move(mask.first)), + mask.second); } else { auto mask = detail::valid_if( input_device_view.begin(), input_device_view.end(), pred, stream, mr); - return std::pair(std::make_unique>(std::move(mask.first)), - mask.second); + return std::pair( + std::make_unique>(std::move(mask.first)), + mask.second); } } } // namespace detail -std::pair>, cudf::size_type> bools_to_mask( +std::pair>, cudf::size_type> bools_to_mask( column_view const& input, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/transform/nans_to_nulls.cu b/cpp/src/transform/nans_to_nulls.cu index 6567d6fc0ff5..b2639500f780 100644 --- a/cpp/src/transform/nans_to_nulls.cu +++ b/cpp/src/transform/nans_to_nulls.cu @@ -24,7 +24,7 @@ namespace cudf { namespace detail { struct dispatch_nan_to_null { template - std::pair>, cudf::size_type> operator()( + std::pair>, cudf::size_type> operator()( column_view const& input, cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(std::is_floating_point_v) { @@ -42,12 +42,13 @@ struct dispatch_nan_to_null { stream, mr); - return std::pair(std::make_unique>(std::move(mask.first)), - mask.second); + return std::pair( + std::make_unique>(std::move(mask.first)), + mask.second); } template - std::pair>, cudf::size_type> operator()( + std::pair>, cudf::size_type> operator()( column_view const& input, cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(!std::is_floating_point_v) { diff --git a/cpp/src/transform/transform.cu b/cpp/src/transform/transform.cu index 9706db3ad80c..7a091139826c 100644 --- a/cpp/src/transform/transform.cu +++ b/cpp/src/transform/transform.cu @@ -59,7 +59,7 @@ struct fixed_width_column { static auto make(data_type type, size_type size, - cuda::device_buffer null_mask, + cuda::device_buffer null_mask, size_type null_count, cuda::stream_ref stream, rmm::device_async_resource_ref mr) @@ -97,12 +97,12 @@ struct mutable_string_views_column_view { struct string_views_column { cuda::device_buffer _data; size_type _size{0}; - cuda::device_buffer _null_mask = + cuda::device_buffer _null_mask = cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); size_type _null_count{0}; static auto make(size_type size, - cuda::device_buffer null_mask, + cuda::device_buffer null_mask, size_type null_count, cuda::stream_ref stream, rmm::device_async_resource_ref mr) @@ -113,19 +113,14 @@ struct string_views_column { auto mutable_view() { - return mutable_string_views_column_view{ - _data.data(), - _size, - reinterpret_cast(_null_mask.data()), - 0, - _null_count}; + return mutable_string_views_column_view{_data.data(), _size, _null_mask.data(), 0, _null_count}; } void set_null_count(size_type count) { _null_count = count; } bool nullable() const { return _null_mask.size() != 0; } - bitmask_type* null_mask() { return reinterpret_cast(_null_mask.data()); } + bitmask_type* null_mask() { return _null_mask.data(); } }; struct mutable_strings_column_view { @@ -143,7 +138,7 @@ struct mutable_strings_column { static auto make(size_type size, rmm::device_buffer chars, std::unique_ptr offsets, - cuda::device_buffer null_mask, + cuda::device_buffer null_mask, size_type null_count) { return mutable_strings_column{make_strings_column( @@ -1000,7 +995,7 @@ rmm::device_uvector make_chars_buffer(column_view const& offsets_view, } std::unique_ptr make_strings_column(device_span strings, - cuda::device_buffer null_mask, + cuda::device_buffer null_mask, size_type null_count, cuda::stream_ref stream, rmm::device_async_resource_ref mr) @@ -1009,7 +1004,7 @@ std::unique_ptr make_strings_column(device_span strin auto size = static_cast(strings.size()); if (size == 0) return make_empty_column(type_id::STRING); - auto stencil = reinterpret_cast(null_mask.data()); + auto stencil = null_mask.data(); // build offsets column from the strings sizes auto sizes = detail::make_counting_transform_iterator( diff --git a/cpp/src/unary/math_ops.cu b/cpp/src/unary/math_ops.cu index 24529ee73d0f..95ed007e5474 100644 --- a/cpp/src/unary/math_ops.cu +++ b/cpp/src/unary/math_ops.cu @@ -336,7 +336,7 @@ std::unique_ptr unary_op_with(column_view const& input, template std::unique_ptr transform_fn(InputIterator begin, InputIterator end, - cuda::device_buffer&& null_mask, + cuda::device_buffer&& null_mask, size_type null_count, cuda::stream_ref stream, rmm::device_async_resource_ref mr) @@ -346,7 +346,7 @@ std::unique_ptr transform_fn(InputIterator begin, std::unique_ptr output = make_fixed_width_column(data_type{type_to_id()}, size, - std::forward>(null_mask), + std::forward>(null_mask), null_count, stream, mr); diff --git a/cpp/tests/bitmask/bitmask_tests.cpp b/cpp/tests/bitmask/bitmask_tests.cpp index 29eff39d992c..6cc0b6573fc4 100644 --- a/cpp/tests/bitmask/bitmask_tests.cpp +++ b/cpp/tests/bitmask/bitmask_tests.cpp @@ -23,7 +23,9 @@ #include +#include #include +#include struct BitmaskUtilitiesTest : public cudf::test::BaseFixture {}; @@ -46,6 +48,36 @@ TEST_F(BitmaskUtilitiesTest, BitmaskAllocationSize) EXPECT_EQ(192u, cudf::bitmask_allocation_size_bytes(1025)); } +TEST_F(BitmaskUtilitiesTest, CreateNullMaskPaddingAndAlignment) +{ + auto const stream = cudf::get_default_stream(); + for (auto const size : {0, 1, 31, 32, 33, 512, 513, 1025}) { + for (auto const state : {cudf::mask_state::UNALLOCATED, + cudf::mask_state::UNINITIALIZED, + cudf::mask_state::ALL_NULL, + cudf::mask_state::ALL_VALID}) { + auto mask = cudf::create_null_mask(size, state, stream); + static_assert(std::is_same_v>); + auto const expected_bytes = + state == cudf::mask_state::UNALLOCATED ? 0 : cudf::bitmask_allocation_size_bytes(size); + EXPECT_EQ(mask.size() * sizeof(cudf::bitmask_type), expected_bytes); + EXPECT_EQ(mask.alignment(), alignof(cudf::bitmask_type)); + if (expected_bytes == 0) { + EXPECT_EQ(mask.data(), nullptr); + continue; + } + EXPECT_EQ(reinterpret_cast(mask.data()) % alignof(cudf::bitmask_type), 0); + if (state == cudf::mask_state::UNINITIALIZED) { continue; } + std::vector words(mask.size()); + CUDF_CUDA_TRY(cudaMemcpyAsync( + words.data(), mask.data(), expected_bytes, cudaMemcpyDefault, stream.get())); + stream.sync(); + auto const expected_word = state == cudf::mask_state::ALL_VALID ? ~cudf::bitmask_type{0} : 0; + EXPECT_THAT(words, ::testing::Each(expected_word)); + } + } +} + TEST_F(BitmaskUtilitiesTest, NumBitmaskWords) { EXPECT_EQ(0, cudf::num_bitmask_words(0)); @@ -387,9 +419,9 @@ struct CopyBitmaskTest : public cudf::test::BaseFixture, cudf::test::UniformRand CopyBitmaskTest() : cudf::test::UniformRandomGenerator{0, 1} {} }; -void cleanEndWord(cuda::device_buffer& mask, int begin_bit, int end_bit) +void cleanEndWord(cuda::device_buffer& mask, int begin_bit, int end_bit) { - auto ptr = reinterpret_cast(mask.data()); + auto ptr = mask.data(); auto number_of_mask_words = cudf::num_bitmask_words(static_cast(end_bit - begin_bit)); auto number_of_bits = end_bit - begin_bit; @@ -442,8 +474,7 @@ TEST_F(CopyBitmaskTest, TestZeroOffset) auto gold_splice_mask = std::get<0>(cudf::test::detail::make_null_mask( validity_bit.begin() + begin_bit, validity_bit.begin() + end_bit)); - auto splice_mask = cudf::copy_bitmask( - reinterpret_cast(input_mask.data()), begin_bit, end_bit); + auto splice_mask = cudf::copy_bitmask(input_mask.data(), begin_bit, end_bit); cleanEndWord(splice_mask, begin_bit, end_bit); auto number_of_bits = end_bit - begin_bit; @@ -465,8 +496,7 @@ TEST_F(CopyBitmaskTest, TestNonZeroOffset) auto gold_splice_mask = std::get<0>(cudf::test::detail::make_null_mask( validity_bit.begin() + begin_bit, validity_bit.begin() + end_bit)); - auto splice_mask = cudf::copy_bitmask( - reinterpret_cast(input_mask.data()), begin_bit, end_bit); + auto splice_mask = cudf::copy_bitmask(input_mask.data(), begin_bit, end_bit); cleanEndWord(splice_mask, begin_bit, end_bit); auto number_of_bits = end_bit - begin_bit; @@ -485,8 +515,7 @@ TEST_F(CopyBitmaskTest, TestCopyColumnViewVectorContiguous) auto [gold_mask, null_count] = cudf::test::detail::make_null_mask(validity_bit.begin(), validity_bit.end()); - auto copy_mask = cudf::copy_bitmask( - reinterpret_cast(gold_mask.data()), 0, num_elements); + auto copy_mask = cudf::copy_bitmask(gold_mask.data(), 0, num_elements); cudf::column original{t, num_elements, rmm::device_buffer{num_elements * sizeof(int), cudf::get_default_stream()}, @@ -510,8 +539,8 @@ TEST_F(CopyBitmaskTest, TestCopyColumnViewVectorContiguous) 760, 760, num_elements}; - std::vector views = cudf::slice(original, indices); - cuda::device_buffer concatenated_bitmask = cudf::concatenate_masks(views); + std::vector views = cudf::slice(original, indices); + cuda::device_buffer concatenated_bitmask = cudf::concatenate_masks(views); cleanEndWord(concatenated_bitmask, 0, num_elements); CUDF_TEST_EXPECT_EQUAL_BUFFERS( concatenated_bitmask.data(), gold_mask.data(), cudf::num_bitmask_words(num_elements)); @@ -542,7 +571,7 @@ TEST_F(CopyBitmaskTest, TestCopyColumnViewVectorDiscontiguous) null_count); views.push_back(cols.back()); } - cuda::device_buffer concatenated_bitmask = cudf::concatenate_masks(views); + cuda::device_buffer concatenated_bitmask = cudf::concatenate_masks(views); cleanEndWord(concatenated_bitmask, 0, num_elements); CUDF_TEST_EXPECT_EQUAL_BUFFERS( concatenated_bitmask.data(), gold_mask.data(), cudf::num_bitmask_words(num_elements)); diff --git a/cpp/tests/bitmask/set_nullmask_tests.cu b/cpp/tests/bitmask/set_nullmask_tests.cu index 052e08c85f36..1fbc60449398 100644 --- a/cpp/tests/bitmask/set_nullmask_tests.cu +++ b/cpp/tests/bitmask/set_nullmask_tests.cu @@ -55,10 +55,11 @@ struct SetBitmaskTest : public cudf::test::BaseFixture { { thrust::host_vector expected(end - begin, valid); // TEST - cuda::device_buffer mask = create_null_mask(size, cudf::mask_state::UNINITIALIZED); + cuda::device_buffer mask = + create_null_mask(size, cudf::mask_state::UNINITIALIZED); // valid ? cudf::mask_state::ALL_NULL : cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(mask.data()), begin, end, valid); - expect_bitmask_equal(reinterpret_cast(mask.data()), begin, expected); + cudf::set_null_mask(mask.data(), begin, end, valid); + expect_bitmask_equal(mask.data(), begin, expected); } void test_null_partition(cudf::size_type size, cudf::size_type middle, bool valid) @@ -69,10 +70,11 @@ struct SetBitmaskTest : public cudf::test::BaseFixture { return (!valid) ^ (i < middle); }); // TEST - cuda::device_buffer mask = create_null_mask(size, cudf::mask_state::UNINITIALIZED); - cudf::set_null_mask(reinterpret_cast(mask.data()), 0, middle, valid); - cudf::set_null_mask(reinterpret_cast(mask.data()), middle, size, !valid); - expect_bitmask_equal(reinterpret_cast(mask.data()), 0, expected); + cuda::device_buffer mask = + create_null_mask(size, cudf::mask_state::UNINITIALIZED); + cudf::set_null_mask(mask.data(), 0, middle, valid); + cudf::set_null_mask(mask.data(), middle, size, !valid); + expect_bitmask_equal(mask.data(), 0, expected); } void test_null_partition_bulk_unsafe(cudf::size_type size, cudf::size_type middle, bool valid) @@ -86,11 +88,12 @@ struct SetBitmaskTest : public cudf::test::BaseFixture { expected2[i] = (valid) ^ (i < middle); }); // TEST - cuda::device_buffer mask1 = create_null_mask(size, cudf::mask_state::UNINITIALIZED); - cuda::device_buffer mask2 = create_null_mask(size, cudf::mask_state::UNINITIALIZED); + cuda::device_buffer mask1 = + create_null_mask(size, cudf::mask_state::UNINITIALIZED); + cuda::device_buffer mask2 = + create_null_mask(size, cudf::mask_state::UNINITIALIZED); - std::vector masks{reinterpret_cast(mask1.data()), - reinterpret_cast(mask2.data())}; + std::vector masks{mask1.data(), mask2.data()}; // Set first halves of bitmasks std::vector begins{0, 0}; @@ -109,8 +112,8 @@ struct SetBitmaskTest : public cudf::test::BaseFixture { masks, begins, ends, cudf::host_span{valids.data(), valids.size()}); // Verify bitmasks - expect_bitmask_equal(reinterpret_cast(mask1.data()), 0, expected1); - expect_bitmask_equal(reinterpret_cast(mask2.data()), 0, expected2); + expect_bitmask_equal(mask1.data(), 0, expected1); + expect_bitmask_equal(mask2.data(), 0, expected2); } void test_null_partition_bulk_safe(cudf::size_type size, cudf::size_type middle, bool valid) @@ -121,19 +124,19 @@ struct SetBitmaskTest : public cudf::test::BaseFixture { return (!valid) ^ (i < middle); }); - cuda::device_buffer mask = create_null_mask(size, cudf::mask_state::UNINITIALIZED); + cuda::device_buffer mask = + create_null_mask(size, cudf::mask_state::UNINITIALIZED); std::vector begins{0, middle}; std::vector ends{middle, size}; - std::vector masks{reinterpret_cast(mask.data()), - reinterpret_cast(mask.data())}; + std::vector masks{mask.data(), mask.data()}; auto valids = cudf::detail::make_host_vector(masks.size(), cudf::get_default_stream()); valids[0] = valid; valids[1] = !valid; cudf::set_null_masks_safe(masks, begins, ends, valids); // Verify bitmasks - expect_bitmask_equal(reinterpret_cast(mask.data()), 0, expected); + expect_bitmask_equal(mask.data(), 0, expected); } }; diff --git a/cpp/tests/column/column_test.cpp b/cpp/tests/column/column_test.cpp index fc8fd1a18ef7..2749f62e8ec2 100644 --- a/cpp/tests/column/column_test.cpp +++ b/cpp/tests/column/column_test.cpp @@ -40,12 +40,15 @@ struct TypedColumnTest : public cudf::test::BaseFixture { : data{_num_elements * sizeof(T), stream}, mask{cudf::create_null_mask(_num_elements, cudf::mask_state::UNINITIALIZED, stream)} { - std::vector h_data(std::max(data.size(), mask.size())); + std::vector h_data(std::max(data.size(), mask.size() * sizeof(cudf::bitmask_type))); std::iota(h_data.begin(), h_data.end(), 0); CUDF_CUDA_TRY( cudaMemcpyAsync(data.data(), h_data.data(), data.size(), cudaMemcpyDefault, stream.get())); - CUDF_CUDA_TRY( - cudaMemcpyAsync(mask.data(), h_data.data(), mask.size(), cudaMemcpyDefault, stream.get())); + CUDF_CUDA_TRY(cudaMemcpyAsync(mask.data(), + h_data.data(), + mask.size() * sizeof(cudf::bitmask_type), + cudaMemcpyDefault, + stream.get())); } cudf::size_type num_elements() { return _num_elements; } @@ -55,10 +58,11 @@ struct TypedColumnTest : public cudf::test::BaseFixture { std::uniform_int_distribution distribution{200, 1000}; cudf::size_type _num_elements{distribution(generator)}; rmm::device_buffer data{}; - cuda::device_buffer mask = cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); - cuda::device_buffer all_valid_mask = + cuda::device_buffer mask = + cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); + cuda::device_buffer all_valid_mask = create_null_mask(num_elements(), cudf::mask_state::ALL_VALID); - cuda::device_buffer all_null_mask = + cuda::device_buffer all_null_mask = create_null_mask(num_elements(), cudf::mask_state::ALL_NULL); }; @@ -190,6 +194,22 @@ TYPED_TEST(TypedColumnTest, SetInvalidSizeNullMaskNonZeroNullCount) cudf::logic_error); } +TYPED_TEST(TypedColumnTest, MoveNullMask) +{ + cudf::column col{this->type(), + this->num_elements(), + std::move(this->data), + cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), + 0}; + auto null_mask = cudf::create_null_mask(this->num_elements(), cudf::mask_state::ALL_NULL); + auto const original_mask = null_mask.data(); + + col.set_null_mask(std::move(null_mask), this->num_elements()); + + EXPECT_EQ(col.view().null_mask(), original_mask); + EXPECT_EQ(col.null_count(), this->num_elements()); +} + TYPED_TEST(TypedColumnTest, SetNullCountEmptyMask) { cudf::column col{this->type(), @@ -281,14 +301,11 @@ TYPED_TEST(TypedColumnTest, MoveDataNoMask) TYPED_TEST(TypedColumnTest, CopyDataAndMask) { - cudf::column col{ - this->type(), - this->num_elements(), - rmm::device_buffer{this->data, cudf::get_default_stream()}, - cudf::copy_bitmask(reinterpret_cast(this->all_valid_mask.data()), - 0, - this->num_elements()), - 0}; + cudf::column col{this->type(), + this->num_elements(), + rmm::device_buffer{this->data, cudf::get_default_stream()}, + cudf::copy_bitmask(this->all_valid_mask.data(), 0, this->num_elements()), + 0}; EXPECT_EQ(this->type(), col.type()); EXPECT_TRUE(col.nullable()); EXPECT_EQ(0, col.null_count()); @@ -300,10 +317,10 @@ TYPED_TEST(TypedColumnTest, CopyDataAndMask) // Verify deep copy cudf::column_view v = col; EXPECT_NE(v.head(), this->data.data()); - EXPECT_NE(v.null_mask(), - reinterpret_cast(this->all_valid_mask.data())); + EXPECT_NE(v.null_mask(), this->all_valid_mask.data()); CUDF_TEST_EXPECT_EQUAL_BUFFERS(v.head(), this->data.data(), this->data.size()); - CUDF_TEST_EXPECT_EQUAL_BUFFERS(v.null_mask(), this->all_valid_mask.data(), this->mask.size()); + CUDF_TEST_EXPECT_EQUAL_BUFFERS( + v.null_mask(), this->all_valid_mask.data(), this->mask.size() * sizeof(cudf::bitmask_type)); } TYPED_TEST(TypedColumnTest, MoveDataAndMask) @@ -423,7 +440,7 @@ TYPED_TEST(TypedColumnTest, DeviceUvectorConstructorWithMask) auto original = cudf::detail::make_device_uvector_async( data, cudf::get_default_stream(), cudf::get_current_device_resource_ref()); auto original_data = original.data(); - auto original_mask = reinterpret_cast(this->all_valid_mask.data()); + auto original_mask = this->all_valid_mask.data(); cudf::column moved_to{std::move(original), std::move(this->all_valid_mask), 0}; verify_column_views(moved_to); @@ -441,27 +458,20 @@ TYPED_TEST(TypedColumnTest, ConstructWithChildren) cudf::data_type{cudf::type_id::INT8}, 42, rmm::device_buffer{this->data, cudf::get_default_stream()}, - cudf::copy_bitmask(reinterpret_cast(this->all_valid_mask.data()), - 0, - this->num_elements()), + cudf::copy_bitmask(this->all_valid_mask.data(), 0, this->num_elements()), 0)); children.emplace_back(std::make_unique( cudf::data_type{cudf::type_id::FLOAT64}, 314, rmm::device_buffer{this->data, cudf::get_default_stream()}, - cudf::copy_bitmask(reinterpret_cast(this->all_valid_mask.data()), - 0, - this->num_elements()), + cudf::copy_bitmask(this->all_valid_mask.data(), 0, this->num_elements()), 0)); - cudf::column col{ - this->type(), - this->num_elements(), - rmm::device_buffer{this->data, cudf::get_default_stream()}, - cudf::copy_bitmask(reinterpret_cast(this->all_valid_mask.data()), - 0, - this->num_elements()), - 0, - std::move(children)}; + cudf::column col{this->type(), + this->num_elements(), + rmm::device_buffer{this->data, cudf::get_default_stream()}, + cudf::copy_bitmask(this->all_valid_mask.data(), 0, this->num_elements()), + 0, + std::move(children)}; verify_column_views(col); EXPECT_EQ(2, col.num_children()); @@ -480,7 +490,7 @@ TYPED_TEST(TypedColumnTest, ReleaseNoChildren) cudf::column::contents contents = col.release(); EXPECT_EQ(original_data, contents.data->data()); - EXPECT_EQ(original_mask, reinterpret_cast(contents.null_mask->data())); + EXPECT_EQ(original_mask, contents.null_mask->data()); EXPECT_EQ(0u, contents.children.size()); EXPECT_EQ(0, col.size()); EXPECT_EQ(0, col.null_count()); @@ -495,34 +505,27 @@ TYPED_TEST(TypedColumnTest, ReleaseWithChildren) this->type(), this->num_elements(), rmm::device_buffer{this->data, cudf::get_default_stream()}, - cudf::copy_bitmask(reinterpret_cast(this->all_valid_mask.data()), - 0, - this->num_elements()), + cudf::copy_bitmask(this->all_valid_mask.data(), 0, this->num_elements()), 0)); children.emplace_back(std::make_unique( this->type(), this->num_elements(), rmm::device_buffer{this->data, cudf::get_default_stream()}, - cudf::copy_bitmask(reinterpret_cast(this->all_valid_mask.data()), - 0, - this->num_elements()), + cudf::copy_bitmask(this->all_valid_mask.data(), 0, this->num_elements()), 0)); - cudf::column col{ - this->type(), - this->num_elements(), - rmm::device_buffer{this->data, cudf::get_default_stream()}, - cudf::copy_bitmask(reinterpret_cast(this->all_valid_mask.data()), - 0, - this->num_elements()), - 0, - std::move(children)}; + cudf::column col{this->type(), + this->num_elements(), + rmm::device_buffer{this->data, cudf::get_default_stream()}, + cudf::copy_bitmask(this->all_valid_mask.data(), 0, this->num_elements()), + 0, + std::move(children)}; auto original_data = col.view().head(); auto original_mask = col.view().null_mask(); cudf::column::contents contents = col.release(); EXPECT_EQ(original_data, contents.data->data()); - EXPECT_EQ(original_mask, reinterpret_cast(contents.null_mask->data())); + EXPECT_EQ(original_mask, contents.null_mask->data()); EXPECT_EQ(2u, contents.children.size()); EXPECT_EQ(0, col.size()); EXPECT_EQ(0, col.null_count()); diff --git a/cpp/tests/column/factories_test.cpp b/cpp/tests/column/factories_test.cpp index b4c727d5b12f..c2c4678999af 100644 --- a/cpp/tests/column/factories_test.cpp +++ b/cpp/tests/column/factories_test.cpp @@ -459,7 +459,7 @@ TYPED_TEST(ListsDictionaryLeafTest, FromNested) DCW leaf({1, 3, -1, 1, 3, 1, 3, -1, 1, 3}, {1, 1, 0, 1, 1, 1, 1, 0, 1, 1}); offset_t offsets{0, 3, 3, 6, 6, 10}; auto mask = cudf::create_null_mask(5, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(mask.data()), 1, 2, false); + cudf::set_null_mask(mask.data(), 1, 2, false); auto data = cudf::make_lists_column(5, offsets.release(), leaf.release(), 0, std::move(mask)); auto s = cudf::make_list_scalar(*data); @@ -471,9 +471,9 @@ TYPED_TEST(ListsDictionaryLeafTest, FromNested) {1, 1, 0, 1, 1, 1, 1, 0, 1, 1, 1, 1, 0, 1, 1, 1, 1, 0, 1, 1, 1, 1, 0, 1, 1, 1, 1, 0, 1, 1}); offset_t offsets2{0, 3, 3, 6, 6, 10, 13, 13, 16, 16, 20, 23, 23, 26, 26, 30}; auto mask2 = cudf::create_null_mask(15, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(mask2.data()), 1, 2, false); - cudf::set_null_mask(reinterpret_cast(mask2.data()), 6, 7, false); - cudf::set_null_mask(reinterpret_cast(mask2.data()), 11, 12, false); + cudf::set_null_mask(mask2.data(), 1, 2, false); + cudf::set_null_mask(mask2.data(), 6, 7, false); + cudf::set_null_mask(mask2.data(), 11, 12, false); auto nested = [&] { auto tmp = cudf::make_lists_column(15, offsets2.release(), leaf2.release(), 3, std::move(mask2)); @@ -587,7 +587,7 @@ TYPED_TEST(ListsStructsLeafTest, FromNested) LCWinner_t({{}, {42}}, valid_t{1, 1}.begin()), valid_t{0, 1}.begin()); auto mask = cudf::create_null_mask(3, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(mask.data()), 0, 1, false); + cudf::set_null_mask(mask.data(), 0, 1, false); auto data = cudf::make_lists_column(3, offset_t{0, 0, 1, 2}.release(), leaf.release(), 1, std::move(mask)); auto s = cudf::make_list_scalar(*data); @@ -601,9 +601,9 @@ TYPED_TEST(ListsStructsLeafTest, FromNested) LCWinner_t({{}, {42}, {}, {42}, {}, {42}}, valid_t{1, 1, 1, 1, 1, 1}.begin()), valid_t{0, 1, 0, 1, 0, 1}.begin()); auto mask2 = cudf::create_null_mask(9, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(mask2.data()), 0, 1, false); - cudf::set_null_mask(reinterpret_cast(mask2.data()), 3, 4, false); - cudf::set_null_mask(reinterpret_cast(mask2.data()), 6, 7, false); + cudf::set_null_mask(mask2.data(), 0, 1, false); + cudf::set_null_mask(mask2.data(), 3, 4, false); + cudf::set_null_mask(mask2.data(), 6, 7, false); auto data2 = [&] { auto tmp = cudf::make_lists_column( 9, offset_t{0, 0, 1, 2, 2, 3, 4, 4, 5, 6}.release(), leaf2.release(), 3, std::move(mask2)); diff --git a/cpp/tests/copying/concatenate_tests.cpp b/cpp/tests/copying/concatenate_tests.cpp index 5a95f8f8069e..f4519d3d4fc4 100644 --- a/cpp/tests/copying/concatenate_tests.cpp +++ b/cpp/tests/copying/concatenate_tests.cpp @@ -54,14 +54,16 @@ struct TypedColumnTest : public cudf::test::BaseFixture { auto typed_mask = mask.data(); std::vector h_data(data.size()); std::iota(h_data.begin(), h_data.end(), char{0}); - std::vector h_mask(mask.size()); + std::vector h_mask(mask.size() * sizeof(cudf::bitmask_type)); std::iota(h_mask.begin(), h_mask.end(), char{0}); CUDF_CUDA_TRY( cudaMemcpyAsync(typed_data, h_data.data(), data.size(), cudaMemcpyDefault, stream.get())); - CUDF_CUDA_TRY( - cudaMemcpyAsync(typed_mask, h_mask.data(), mask.size(), cudaMemcpyDefault, stream.get())); - _null_count = - cudf::null_count(reinterpret_cast(mask.data()), 0, _num_elements); + CUDF_CUDA_TRY(cudaMemcpyAsync(typed_mask, + h_mask.data(), + mask.size() * sizeof(cudf::bitmask_type), + cudaMemcpyDefault, + stream.get())); + _null_count = cudf::null_count(mask.data(), 0, _num_elements); stream.sync(); } @@ -73,11 +75,12 @@ struct TypedColumnTest : public cudf::test::BaseFixture { std::uniform_int_distribution distribution{200, 1000}; cudf::size_type _num_elements{distribution(generator)}; rmm::device_buffer data{}; - cuda::device_buffer mask = cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); + cuda::device_buffer mask = + cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); cudf::size_type _null_count{}; - cuda::device_buffer all_valid_mask = + cuda::device_buffer all_valid_mask = create_null_mask(num_elements(), cudf::mask_state::ALL_VALID); - cuda::device_buffer all_null_mask = + cuda::device_buffer all_null_mask = create_null_mask(num_elements(), cudf::mask_state::ALL_NULL); }; diff --git a/cpp/tests/copying/get_value_tests.cpp b/cpp/tests/copying/get_value_tests.cpp index 82123c0e9807..bf7e5ec8424e 100644 --- a/cpp/tests/copying/get_value_tests.cpp +++ b/cpp/tests/copying/get_value_tests.cpp @@ -506,15 +506,13 @@ struct ListGetStructValueTest : public cudf::test::BaseFixture { auto d_null_mask = cudf::create_null_mask( num_lists, null_count == 0 ? cudf::mask_state::UNALLOCATED : cudf::mask_state::ALL_NULL); if (null_count > 0) { - std::for_each( - cuda::counting_iterator{0}, - cuda::counting_iterator{num_lists}, - [&](auto i) { - if (*(null_mask.begin() + i)) { - cudf::set_null_mask( - reinterpret_cast(d_null_mask.data()), i, i + 1, true); - } - }); + std::for_each(cuda::counting_iterator{0}, + cuda::counting_iterator{num_lists}, + [&](auto i) { + if (*(null_mask.begin() + i)) { + cudf::set_null_mask(d_null_mask.data(), i, i + 1, true); + } + }); } return cudf::make_lists_column( num_lists, offsets.release(), std::move(child), null_count, std::move(d_null_mask)); diff --git a/cpp/tests/copying/scatter_list_scalar_tests.cpp b/cpp/tests/copying/scatter_list_scalar_tests.cpp index afd9d2d09c48..da73f602b5ca 100644 --- a/cpp/tests/copying/scatter_list_scalar_tests.cpp +++ b/cpp/tests/copying/scatter_list_scalar_tests.cpp @@ -386,7 +386,7 @@ TYPED_TEST(ScatterListOfStructScalarTest, NullScalar) offset_t ex_offsets{0, 0, 0, 1, 1}; auto null_mask = cudf::create_null_mask(4, cudf::mask_state::ALL_NULL); - cudf::set_null_mask(reinterpret_cast(null_mask.data()), 2, 3, true); + cudf::set_null_mask(null_mask.data(), 2, 3, true); auto expected = cudf::make_lists_column(4, ex_offsets.release(), ex_child.release(), 3, std::move(null_mask)); @@ -412,7 +412,7 @@ TYPED_TEST(ScatterListOfStructScalarTest, NullableTargetRow) {1, 1, 1, 0, 1, 1}); offset_t offsets{0, 2, 2, 3, 6}; auto null_mask = cudf::create_null_mask(4, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(null_mask.data()), 1, 3, false); + cudf::set_null_mask(null_mask.data(), 1, 3, false); auto col = cudf::make_lists_column(4, offsets.release(), child.release(), 2, std::move(null_mask)); @@ -428,7 +428,7 @@ TYPED_TEST(ScatterListOfStructScalarTest, NullableTargetRow) offset_t ex_offsets{0, 2, 2, 5, 8}; auto ex_null_mask = cudf::create_null_mask(4, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(ex_null_mask.data()), 1, 2, false); + cudf::set_null_mask(ex_null_mask.data(), 1, 2, false); auto expected = cudf::make_lists_column( 4, ex_offsets.release(), ex_child.release(), 1, std::move(ex_null_mask)); diff --git a/cpp/tests/dictionary/factories_test.cpp b/cpp/tests/dictionary/factories_test.cpp index 9cc69b8fb0ad..75cb76bce682 100644 --- a/cpp/tests/dictionary/factories_test.cpp +++ b/cpp/tests/dictionary/factories_test.cpp @@ -66,8 +66,9 @@ TEST_F(DictionaryFactoriesTest, ColumnsWithNulls) cudf::test::fixed_width_column_wrapper keys(h_keys.begin(), h_keys.end()); std::vector h_values{1, 2, 3, 1, 2, 3, 0}; cudf::test::fixed_width_column_wrapper values(h_values.begin(), h_values.end()); - auto size = static_cast(h_values.size()); - cuda::device_buffer null_mask = create_null_mask(size, cudf::mask_state::ALL_NULL); + auto size = static_cast(h_values.size()); + cuda::device_buffer null_mask = + create_null_mask(size, cudf::mask_state::ALL_NULL); auto dictionary = cudf::make_dictionary_column(keys.release(), values.release(), std::move(null_mask), size); cudf::dictionary_column_view view(dictionary->view()); diff --git a/cpp/tests/io/experimental/hybrid_scan_common.cpp b/cpp/tests/io/experimental/hybrid_scan_common.cpp index 3c88e99d1c54..87938388c0ac 100644 --- a/cpp/tests/io/experimental/hybrid_scan_common.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_common.cpp @@ -363,12 +363,13 @@ std::pair, std::vector> create_parquet_with_s auto const make_null_mask = [stream](auto begin, auto end) { auto [null_mask, null_count] = cudf::test::detail::make_null_mask_vector(begin, end); - auto const* data = reinterpret_cast(null_mask.data()); - auto d_mask = cuda::device_buffer{ + auto const* data = null_mask.data(); + auto d_mask = cuda::device_buffer{ stream, cudf::get_current_device_resource_ref(), data, - data + cudf::bitmask_allocation_size_bytes(cudf::distance(begin, end))}; + data + cudf::bitmask_allocation_size_bytes(cudf::distance(begin, end)) / + sizeof(cudf::bitmask_type)}; stream.sync(); return std::pair{std::move(d_mask), null_count}; }; diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 3d5f0b254454..25c74f623aa3 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -2773,7 +2773,7 @@ TEST_F(CastVariantStatusTest, SqlNullInputProducesRowNullStatus) cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)); // Mask row 1 SQL null auto null_mask = cudf::create_null_mask(2, cudf::mask_state::ALL_VALID, stream, cmr()); - cudf::set_null_mask(reinterpret_cast(null_mask.data()), 1, 2, false); + cudf::set_null_mask(null_mask.data(), 1, 2, false); stream.sync(); values_col->set_null_mask(std::move(null_mask), 1); diff --git a/cpp/tests/io/orc_chunked_reader_test.cu b/cpp/tests/io/orc_chunked_reader_test.cu index 257dbc698332..e46bd1cbaa92 100644 --- a/cpp/tests/io/orc_chunked_reader_test.cu +++ b/cpp/tests/io/orc_chunked_reader_test.cu @@ -71,7 +71,7 @@ auto write_file(std::vector>& input_columns, auto const [null_mask, null_count] = cudf::test::detail::make_null_mask(valid_iter + offset, valid_iter + col->size() + offset); col = cudf::structs::detail::superimpose_and_sanitize_nulls( - reinterpret_cast(null_mask.data()), + null_mask.data(), null_count, std::move(col), cudf::get_default_stream(), diff --git a/cpp/tests/io/parquet_chunked_reader_test.cu b/cpp/tests/io/parquet_chunked_reader_test.cu index b802c9107909..cb415a0da833 100644 --- a/cpp/tests/io/parquet_chunked_reader_test.cu +++ b/cpp/tests/io/parquet_chunked_reader_test.cu @@ -74,7 +74,7 @@ auto write_file(std::vector>& input_columns, auto const [null_mask, null_count] = cudf::test::detail::make_null_mask(valid_iter + offset, valid_iter + col->size() + offset); col = cudf::structs::detail::superimpose_and_sanitize_nulls( - reinterpret_cast(null_mask.data()), + null_mask.data(), null_count, std::move(col), cudf::get_default_stream(), diff --git a/cpp/tests/io/parquet_reader_test.cpp b/cpp/tests/io/parquet_reader_test.cpp index 5510c877ff4a..8016584a9b62 100644 --- a/cpp/tests/io/parquet_reader_test.cpp +++ b/cpp/tests/io/parquet_reader_test.cpp @@ -3742,7 +3742,7 @@ TEST_F(ParquetReaderTest, RepeatedNoAnnotations) auto num_list_rows = list_offsets_column->size() - 1; auto mask = cudf::create_null_mask(6, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(mask.data()), 0, 2, false); + cudf::set_null_mask(mask.data(), 0, 2, false); auto list_col = cudf::make_lists_column( num_list_rows, std::move(list_offsets_column), struct_col.release(), 2, std::move(mask)); diff --git a/cpp/tests/lists/reverse_tests.cpp b/cpp/tests/lists/reverse_tests.cpp index 52b5da95364f..e1eee5f0490c 100644 --- a/cpp/tests/lists/reverse_tests.cpp +++ b/cpp/tests/lists/reverse_tests.cpp @@ -221,7 +221,7 @@ TYPED_TEST(ListsReverseTypedTest, InputListsOfListsWithNulls) .release(); auto offsets = ints_col{0, 0, 3, 3, 6, 9, 10, 10, 10}.release(); auto null_mask = cudf::create_null_mask(8, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(null_mask.data()), 2, 3, false); + cudf::set_null_mask(null_mask.data(), 2, 3, false); return cudf::make_lists_column( 8, std::move(offsets), std::move(child), 1, std::move(null_mask)); @@ -243,7 +243,7 @@ TYPED_TEST(ListsReverseTypedTest, InputListsOfListsWithNulls) .release(); auto offsets = ints_col{0, 0, 3, 3, 6, 9, 10, 10, 10}.release(); auto null_mask = cudf::create_null_mask(8, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(null_mask.data()), 2, 3, false); + cudf::set_null_mask(null_mask.data(), 2, 3, false); return cudf::make_lists_column( 8, std::move(offsets), std::move(child), 1, std::move(null_mask)); @@ -274,7 +274,7 @@ TYPED_TEST(ListsReverseTypedTest, InputListsOfListsWithNulls) .release(); auto offsets = ints_col{0, 3, 3, 6}.release(); auto null_mask = cudf::create_null_mask(3, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(null_mask.data()), 1, 2, false); + cudf::set_null_mask(null_mask.data(), 1, 2, false); return cudf::make_lists_column( 3, std::move(offsets), std::move(child), 1, std::move(null_mask)); @@ -289,7 +289,7 @@ TYPED_TEST(ListsReverseTypedTest, InputListsOfListsWithNulls) auto child = lists_col{{{4, 5, 6}, {} /*null*/, {4, 5}}, null_at(1)}.release(); auto offsets = ints_col{0, 0, 3}.release(); auto null_mask = cudf::create_null_mask(2, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(null_mask.data()), 0, 1, false); + cudf::set_null_mask(null_mask.data(), 0, 1, false); return cudf::make_lists_column( 2, std::move(offsets), std::move(child), 1, std::move(null_mask)); @@ -347,8 +347,8 @@ TYPED_TEST(ListsReverseTypedTest, InputListsOfStructsWithNulls) }(); auto offsets = ints_col{0, 0, 8, 16, 16, 16, 24}.release(); auto null_mask = cudf::create_null_mask(6, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(null_mask.data()), 0, 1, false); - cudf::set_null_mask(reinterpret_cast(null_mask.data()), 4, 5, false); + cudf::set_null_mask(null_mask.data(), 0, 1, false); + cudf::set_null_mask(null_mask.data(), 4, 5, false); return cudf::make_lists_column( 6, std::move(offsets), std::move(child), 2, std::move(null_mask)); @@ -397,8 +397,8 @@ TYPED_TEST(ListsReverseTypedTest, InputListsOfStructsWithNulls) }(); auto offsets = ints_col{0, 0, 8, 16, 16, 16, 24}.release(); auto null_mask = cudf::create_null_mask(6, cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(null_mask.data()), 0, 1, false); - cudf::set_null_mask(reinterpret_cast(null_mask.data()), 4, 5, false); + cudf::set_null_mask(null_mask.data(), 0, 1, false); + cudf::set_null_mask(null_mask.data(), 4, 5, false); return cudf::make_lists_column( 6, std::move(offsets), std::move(child), 2, std::move(null_mask)); diff --git a/cpp/tests/quantiles/percentile_approx_test.cpp b/cpp/tests/quantiles/percentile_approx_test.cpp index 81161d9aaacd..2b5fc00927d5 100644 --- a/cpp/tests/quantiles/percentile_approx_test.cpp +++ b/cpp/tests/quantiles/percentile_approx_test.cpp @@ -227,7 +227,7 @@ void grouped_test(cudf::data_type input_type, std::vector> p }); } -std::pair, cudf::size_type> make_null_mask( +std::pair, cudf::size_type> make_null_mask( cudf::column_view const& col) { auto itr = cudf::test::iterators::valids_at_multiples_of(2); diff --git a/cpp/tests/reductions/list_rank_test.cpp b/cpp/tests/reductions/list_rank_test.cpp index e927a994cfe3..c9735b5c77d5 100644 --- a/cpp/tests/reductions/list_rank_test.cpp +++ b/cpp/tests/reductions/list_rank_test.cpp @@ -184,7 +184,7 @@ TEST_F(ListRankScanTest, ListOfStruct) auto list_column = cudf::column_view(cudf::data_type(cudf::type_id::LIST), 17, nullptr, - reinterpret_cast(null_mask.data()), + null_mask.data(), null_count, 0, {offsets, struct_col}); diff --git a/cpp/tests/rolling/grouped_rolling_range_test.cpp b/cpp/tests/rolling/grouped_rolling_range_test.cpp index b65b4d42ee34..49bd432d8a76 100644 --- a/cpp/tests/rolling/grouped_rolling_range_test.cpp +++ b/cpp/tests/rolling/grouped_rolling_range_test.cpp @@ -160,17 +160,11 @@ struct GroupedRollingRangeOrderByNumericTest : public BaseGroupedRollingRangeOrd auto const nulled_order_by = [&] { auto col = generate_order_by_column(); auto new_null_mask = create_null_mask(col->size(), cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(new_null_mask.data()), - 0, - 2, + cudf::set_null_mask(new_null_mask.data(), 0, 2, false); // Nulls in first group. - cudf::set_null_mask(reinterpret_cast(new_null_mask.data()), - 6, - 8, + cudf::set_null_mask(new_null_mask.data(), 6, 8, false); // Nulls in second group. - cudf::set_null_mask(reinterpret_cast(new_null_mask.data()), - 10, - 12, + cudf::set_null_mask(new_null_mask.data(), 10, 12, false); // Nulls in third group. col->set_null_mask(std::move(new_null_mask), 6); return col; @@ -473,17 +467,11 @@ struct GroupedRollingRangeOrderByDecimalTypedTest auto const nulled_order_by = [&] { auto col = generate_order_by_column(order_by_column_scale); auto new_null_mask = create_null_mask(col->size(), cudf::mask_state::ALL_VALID); - cudf::set_null_mask(reinterpret_cast(new_null_mask.data()), - 0, - 2, + cudf::set_null_mask(new_null_mask.data(), 0, 2, false); // Nulls in first group. - cudf::set_null_mask(reinterpret_cast(new_null_mask.data()), - 6, - 8, + cudf::set_null_mask(new_null_mask.data(), 6, 8, false); // Nulls in second group. - cudf::set_null_mask(reinterpret_cast(new_null_mask.data()), - 10, - 12, + cudf::set_null_mask(new_null_mask.data(), 10, 12, false); // Nulls in third group. col->set_null_mask(std::move(new_null_mask), 6); return col; diff --git a/cpp/tests/stream_compaction/distinct_tests.cpp b/cpp/tests/stream_compaction/distinct_tests.cpp index 4887a69676b0..1212bbb5e720 100644 --- a/cpp/tests/stream_compaction/distinct_tests.cpp +++ b/cpp/tests/stream_compaction/distinct_tests.cpp @@ -872,7 +872,7 @@ TEST_F(DistinctKeepAny, ListsOfStructs) auto const keys = cudf::column_view(cudf::data_type(cudf::type_id::LIST), 17, nullptr, - reinterpret_cast(null_mask.data()), + null_mask.data(), null_count, 0, {offsets, structs}); @@ -958,7 +958,7 @@ TEST_F(DistinctKeepFirstLastNone, ListsOfStructs) auto const keys = cudf::column_view(cudf::data_type(cudf::type_id::LIST), 17, nullptr, - reinterpret_cast(null_mask.data()), + null_mask.data(), null_count, 0, {offsets, structs}); @@ -1054,7 +1054,7 @@ TEST_F(DistinctKeepAny, SlicedListsOfStructs) auto const keys = cudf::column_view(cudf::data_type(cudf::type_id::LIST), 17, nullptr, - reinterpret_cast(null_mask.data()), + null_mask.data(), null_count, 0, {offsets, structs}); @@ -1104,25 +1104,23 @@ TEST_F(DistinctKeepAny, ListsOfEmptyStructs) auto const structs_null_it = nulls_at({0, 1, 2, 3, 4, 5, 6, 7}); auto [structs_null_mask, structs_null_count] = cudf::test::detail::make_null_mask(structs_null_it, structs_null_it + 14); - auto const structs = - cudf::column_view(cudf::data_type(cudf::type_id::STRUCT), - 14, - nullptr, - reinterpret_cast(structs_null_mask.data()), - structs_null_count); + auto const structs = cudf::column_view(cudf::data_type(cudf::type_id::STRUCT), + 14, + nullptr, + structs_null_mask.data(), + structs_null_count); auto const offsets = int32s_col{0, 0, 0, 0, 0, 2, 4, 6, 7, 8, 9, 10, 12, 14}; auto const lists_null_it = nulls_at({2, 3}); auto [lists_null_mask, lists_null_count] = cudf::test::detail::make_null_mask(lists_null_it, lists_null_it + 13); - auto const keys = - cudf::column_view(cudf::data_type(cudf::type_id::LIST), - 13, - nullptr, - reinterpret_cast(lists_null_mask.data()), - lists_null_count, - 0, - {offsets, structs}); + auto const keys = cudf::column_view(cudf::data_type(cudf::type_id::LIST), + 13, + nullptr, + lists_null_mask.data(), + lists_null_count, + 0, + {offsets, structs}); auto const idx = int32s_col{1, 1, 2, 2, 3, 3, 3, 4, 4, 5, 5, 6, 6}; auto const input = cudf::table_view{{idx, keys}}; diff --git a/cpp/tests/stream_compaction/stable_distinct_tests.cpp b/cpp/tests/stream_compaction/stable_distinct_tests.cpp index 65c9896b6f54..ad95304bd8f2 100644 --- a/cpp/tests/stream_compaction/stable_distinct_tests.cpp +++ b/cpp/tests/stream_compaction/stable_distinct_tests.cpp @@ -812,7 +812,7 @@ TEST_F(StableDistinctKeepAny, ListsOfStructs) auto const keys = cudf::column_view(cudf::data_type(cudf::type_id::LIST), 17, nullptr, - reinterpret_cast(null_mask.data()), + null_mask.data(), null_count, 0, {offsets, structs}); @@ -896,7 +896,7 @@ TEST_F(StableDistinctKeepFirstLastNone, ListsOfStructs) auto const keys = cudf::column_view(cudf::data_type(cudf::type_id::LIST), 17, nullptr, - reinterpret_cast(null_mask.data()), + null_mask.data(), null_count, 0, {offsets, structs}); @@ -989,7 +989,7 @@ TEST_F(StableDistinctKeepAny, SlicedListsOfStructs) auto const keys = cudf::column_view(cudf::data_type(cudf::type_id::LIST), 17, nullptr, - reinterpret_cast(null_mask.data()), + null_mask.data(), null_count, 0, {offsets, structs}); @@ -1040,25 +1040,23 @@ TEST_F(StableDistinctKeepAny, ListsOfEmptyStructs) auto const structs_null_it = nulls_at({0, 1, 2, 3, 4, 5, 6, 7}); auto [structs_null_mask, structs_null_count] = cudf::test::detail::make_null_mask(structs_null_it, structs_null_it + 14); - auto const structs = - cudf::column_view(cudf::data_type(cudf::type_id::STRUCT), - 14, - nullptr, - reinterpret_cast(structs_null_mask.data()), - structs_null_count); + auto const structs = cudf::column_view(cudf::data_type(cudf::type_id::STRUCT), + 14, + nullptr, + structs_null_mask.data(), + structs_null_count); auto const offsets = int32s_col{0, 0, 0, 0, 0, 2, 4, 6, 7, 8, 9, 10, 12, 14}; auto const lists_null_it = nulls_at({2, 3}); auto [lists_null_mask, lists_null_count] = cudf::test::detail::make_null_mask(lists_null_it, lists_null_it + 13); - auto const keys = - cudf::column_view(cudf::data_type(cudf::type_id::LIST), - 13, - nullptr, - reinterpret_cast(lists_null_mask.data()), - lists_null_count, - 0, - {offsets, structs}); + auto const keys = cudf::column_view(cudf::data_type(cudf::type_id::LIST), + 13, + nullptr, + lists_null_mask.data(), + lists_null_count, + 0, + {offsets, structs}); auto const idx = int32s_col{1, 1, 2, 2, 3, 3, 3, 4, 4, 5, 5, 6, 6}; auto const input = cudf::table_view{{idx, keys}}; diff --git a/cpp/tests/stream_compaction/unique_tests.cpp b/cpp/tests/stream_compaction/unique_tests.cpp index 313b63b127b9..f11e4794f04b 100644 --- a/cpp/tests/stream_compaction/unique_tests.cpp +++ b/cpp/tests/stream_compaction/unique_tests.cpp @@ -467,7 +467,7 @@ TEST_F(Unique, ListsOfStructsKeepAny) auto const keys = cudf::column_view(cudf::data_type(cudf::type_id::LIST), 17, nullptr, - reinterpret_cast(null_mask.data()), + null_mask.data(), null_count, 0, {offsets, structs}); @@ -550,7 +550,7 @@ TEST_F(Unique, ListsOfStructsKeepFirstLastNone) auto const keys = cudf::column_view(cudf::data_type(cudf::type_id::LIST), 17, nullptr, - reinterpret_cast(null_mask.data()), + null_mask.data(), null_count, 0, {offsets, structs}); @@ -606,25 +606,23 @@ TEST_F(Unique, ListsOfEmptyStructsKeepAny) auto const structs_null_it = nulls_at({0, 1, 2, 3, 4, 5, 6, 7}); auto [structs_null_mask, structs_null_count] = cudf::test::detail::make_null_mask(structs_null_it, structs_null_it + 14); - auto const structs = - cudf::column_view(cudf::data_type(cudf::type_id::STRUCT), - 14, - nullptr, - reinterpret_cast(structs_null_mask.data()), - structs_null_count); + auto const structs = cudf::column_view(cudf::data_type(cudf::type_id::STRUCT), + 14, + nullptr, + structs_null_mask.data(), + structs_null_count); auto const offsets = int32s_col{0, 0, 0, 0, 0, 2, 4, 6, 7, 8, 9, 10, 12, 14}; auto const lists_null_it = nulls_at({2, 3}); auto [lists_null_mask, lists_null_count] = cudf::test::detail::make_null_mask(lists_null_it, lists_null_it + 13); - auto const keys = - cudf::column_view(cudf::data_type(cudf::type_id::LIST), - 13, - nullptr, - reinterpret_cast(lists_null_mask.data()), - lists_null_count, - 0, - {offsets, structs}); + auto const keys = cudf::column_view(cudf::data_type(cudf::type_id::LIST), + 13, + nullptr, + lists_null_mask.data(), + lists_null_count, + 0, + {offsets, structs}); auto const idx = int32s_col{1, 1, 2, 2, 3, 3, 3, 4, 4, 5, 5, 6, 6}; auto const input = cudf::table_view{{idx, keys}}; diff --git a/cpp/tests/streams/column_view_test.cpp b/cpp/tests/streams/column_view_test.cpp index 0cda5671e23f..bbfc1a3e7fb0 100644 --- a/cpp/tests/streams/column_view_test.cpp +++ b/cpp/tests/streams/column_view_test.cpp @@ -24,12 +24,15 @@ struct TypedColumnTest : public cudf::test::BaseFixture { : data{_num_elements * sizeof(T), stream}, mask{cudf::create_null_mask(_num_elements, cudf::mask_state::UNINITIALIZED, stream)} { - std::vector h_data(std::max(data.size(), mask.size())); + std::vector h_data(std::max(data.size(), mask.size() * sizeof(cudf::bitmask_type))); std::iota(h_data.begin(), h_data.end(), 0); CUDF_CUDA_TRY( cudaMemcpyAsync(data.data(), h_data.data(), data.size(), cudaMemcpyDefault, stream.get())); - CUDF_CUDA_TRY( - cudaMemcpyAsync(mask.data(), h_data.data(), mask.size(), cudaMemcpyDefault, stream.get())); + CUDF_CUDA_TRY(cudaMemcpyAsync(mask.data(), + h_data.data(), + mask.size() * sizeof(cudf::bitmask_type), + cudaMemcpyDefault, + stream.get())); } cudf::size_type num_elements() { return _num_elements; } @@ -39,10 +42,11 @@ struct TypedColumnTest : public cudf::test::BaseFixture { std::uniform_int_distribution distribution{200, 1000}; cudf::size_type _num_elements{distribution(generator)}; rmm::device_buffer data{}; - cuda::device_buffer mask = cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); - cuda::device_buffer all_valid_mask = + cuda::device_buffer mask = + cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED); + cuda::device_buffer all_valid_mask = create_null_mask(num_elements(), cudf::mask_state::ALL_VALID, cudf::test::get_default_stream()); - cuda::device_buffer all_null_mask = + cuda::device_buffer all_null_mask = create_null_mask(num_elements(), cudf::mask_state::ALL_NULL, cudf::test::get_default_stream()); }; diff --git a/cpp/tests/transform/mask_to_bools_test.cpp b/cpp/tests/transform/mask_to_bools_test.cpp index e580ab68f797..abf71590d748 100644 --- a/cpp/tests/transform/mask_to_bools_test.cpp +++ b/cpp/tests/transform/mask_to_bools_test.cpp @@ -53,8 +53,7 @@ TEST_P(MaskToBoolsTest, LargeDataSizeTest) auto mask = cudf::bools_to_mask(col); - auto out = cudf::mask_to_bools( - reinterpret_cast(mask.first->data()), begin_bit, end_bit); + auto out = cudf::mask_to_bools(mask.first->data(), begin_bit, end_bit); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, out->view()); } diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index 41c99b6d7cbd..1aff59475894 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -967,16 +967,16 @@ std::vector bitmask_to_host(cudf::column_view const& c, auto num_bitmasks = num_bitmask_words(c.size()); auto [bitmask_span, _] = [&] { if (c.offset() == 0) { - return std::pair, cuda::device_buffer>{ + return std::pair, + cuda::device_buffer>{ cudf::device_span(c.null_mask(), num_bitmasks), cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)}; } auto mask = copy_bitmask( c.null_mask(), c.offset(), c.offset() + c.size(), stream, mr.get_temporary_mr()); - return std::pair, cuda::device_buffer>{ - cudf::device_span(reinterpret_cast(mask.data()), - num_bitmasks), - std::move(mask)}; + return std::pair, + cuda::device_buffer>{ + cudf::device_span(mask.data(), num_bitmasks), std::move(mask)}; }(); return cudf::detail::make_std_vector(bitmask_span, stream); } else { diff --git a/java/src/main/native/src/ColumnViewJni.cpp b/java/src/main/native/src/ColumnViewJni.cpp index 0931b16cb54d..e25cb365f796 100644 --- a/java/src/main/native/src/ColumnViewJni.cpp +++ b/java/src/main/native/src/ColumnViewJni.cpp @@ -2248,7 +2248,7 @@ JNIEXPORT jlong JNICALL Java_ai_rapids_cudf_ColumnView_bitwiseMergeAndSetValidit // descendants for STRUCTs so that child masks stay consistent , // and fix offsets for LIST/STRINGs by purging non-empty nulls. auto result = cudf::structs::detail::superimpose_and_sanitize_nulls( - reinterpret_cast(merge_mask.data()), + merge_mask.data(), merge_null_count, std::move(copy), cudf::get_default_stream(), diff --git a/java/src/main/native/src/ColumnViewJni.cu b/java/src/main/native/src/ColumnViewJni.cu index d75885977e4d..80168c041577 100644 --- a/java/src/main/native/src/ColumnViewJni.cu +++ b/java/src/main/native/src/ColumnViewJni.cu @@ -128,8 +128,7 @@ void post_process_list_overlap(cudf::column_view const& lhs, // new nullmask. if (overlap_cv.nullable()) { auto [null_mask, null_count] = cudf::bitmask_and( - std::vector{ - overlap_cv.null_mask(), reinterpret_cast(new_null_mask->data())}, + std::vector{overlap_cv.null_mask(), new_null_mask->data()}, std::vector{0, 0}, overlap_cv.size(), stream, diff --git a/python/pylibcudf/pylibcudf/gpumemoryview.pxd b/python/pylibcudf/pylibcudf/gpumemoryview.pxd index 4992ec5498e6..7395950dbd4b 100644 --- a/python/pylibcudf/pylibcudf/gpumemoryview.pxd +++ b/python/pylibcudf/pylibcudf/gpumemoryview.pxd @@ -3,17 +3,18 @@ from libc.stdint cimport uint64_t, uintptr_t from libcpp.memory cimport unique_ptr -from pylibcudf.libcudf.utilities.device_buffer cimport byte, device_buffer +from pylibcudf.libcudf.types cimport bitmask_type +from pylibcudf.libcudf.utilities.device_buffer cimport device_buffer cdef class _CudaDeviceBuffer: - cdef unique_ptr[device_buffer[byte]] c_obj + cdef unique_ptr[device_buffer[bitmask_type]] c_obj cdef object stream cdef object mr cdef gpumemoryview _from_cuda_device_buffer( - unique_ptr[device_buffer[byte]] buf, object stream, object mr + unique_ptr[device_buffer[bitmask_type]] buf, object stream, object mr ) cdef class gpumemoryview: diff --git a/python/pylibcudf/pylibcudf/gpumemoryview.pyx b/python/pylibcudf/pylibcudf/gpumemoryview.pyx index aaa50aaf746f..f6048bf2fe53 100644 --- a/python/pylibcudf/pylibcudf/gpumemoryview.pyx +++ b/python/pylibcudf/pylibcudf/gpumemoryview.pyx @@ -1,6 +1,7 @@ # SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 +from pylibcudf.libcudf.types cimport bitmask_type from libc.stddef cimport size_t from libc.stdint cimport uintptr_t, uint64_t from libcpp.memory cimport unique_ptr @@ -12,7 +13,7 @@ import operator from typing import Any from .types cimport DataType, size_of, type_id -from pylibcudf.libcudf.utilities.device_buffer cimport byte, device_buffer +from pylibcudf.libcudf.utilities.device_buffer cimport device_buffer cdef class _CudaDeviceBuffer: @@ -20,14 +21,14 @@ cdef class _CudaDeviceBuffer: def __cuda_array_interface__(self): return { "data": (self.c_obj.get().data(), False), - "shape": (self.c_obj.get().size(),), + "shape": (self.c_obj.get().size() * sizeof(bitmask_type),), "typestr": "|u1", "version": 3, } cdef gpumemoryview _from_cuda_device_buffer( - unique_ptr[device_buffer[byte]] buf, object stream, object mr + unique_ptr[device_buffer[bitmask_type]] buf, object stream, object mr ): cdef _CudaDeviceBuffer owner = _CudaDeviceBuffer.__new__(_CudaDeviceBuffer) owner.c_obj = move(buf) diff --git a/python/pylibcudf/pylibcudf/join.pyx b/python/pylibcudf/pylibcudf/join.pyx index 93425b1a6bfb..b613b4585683 100644 --- a/python/pylibcudf/pylibcudf/join.pyx +++ b/python/pylibcudf/pylibcudf/join.pyx @@ -1,6 +1,7 @@ # SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 +from pylibcudf.libcudf.types cimport bitmask_type from cython.operator import dereference from libc.stddef cimport size_t @@ -13,7 +14,7 @@ from pylibcudf.libcudf.column.column cimport column from pylibcudf.libcudf.table.table cimport table from pylibcudf.libcudf.table.table_view cimport table_view from pylibcudf.libcudf.types cimport mask_state, null_equality -from pylibcudf.libcudf.utilities.device_buffer cimport byte, device_buffer +from pylibcudf.libcudf.utilities.device_buffer cimport device_buffer from rmm.pylibrmm.stream cimport Stream from rmm.pylibrmm.memory_resource cimport DeviceMemoryResource @@ -54,7 +55,7 @@ cdef Column _column_from_gather_map( ): # helper to convert a gather map to a Column cdef Stream _stream = _get_stream(stream) - cdef unique_ptr[device_buffer[byte]] mask = ( + cdef unique_ptr[device_buffer[bitmask_type]] mask = ( cpp_null_mask.create_null_mask_unique_ptr( 0, mask_state.UNALLOCATED, diff --git a/python/pylibcudf/pylibcudf/libcudf/column/column.pxd b/python/pylibcudf/pylibcudf/libcudf/column/column.pxd index 5252da219edf..19bd36b2473e 100644 --- a/python/pylibcudf/pylibcudf/libcudf/column/column.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/column/column.pxd @@ -10,10 +10,10 @@ from pylibcudf.libcudf.column.column_view cimport ( ) from pylibcudf.libcudf.types cimport data_type, size_type from pylibcudf.libcudf.utilities.device_buffer cimport ( - byte, device_buffer as cuda_device_buffer, ) +from pylibcudf.libcudf.types cimport bitmask_type from rmm.librmm.device_buffer cimport device_buffer from cuda.bindings.cyruntime cimport cudaStream_t from rmm.librmm.memory_resource cimport device_async_resource_ref @@ -22,7 +22,7 @@ from rmm.librmm.memory_resource cimport device_async_resource_ref cdef extern from "cudf/column/column.hpp" namespace "cudf" nogil: cdef cppclass column_contents "cudf::column::contents": unique_ptr[device_buffer] data - unique_ptr[cuda_device_buffer[byte]] null_mask + unique_ptr[cuda_device_buffer[bitmask_type]] null_mask vector[unique_ptr[column]] children cdef cppclass column: diff --git a/python/pylibcudf/pylibcudf/libcudf/column/column_factories.pxd b/python/pylibcudf/pylibcudf/libcudf/column/column_factories.pxd index 690f250e8a39..d4fd24c65750 100644 --- a/python/pylibcudf/pylibcudf/libcudf/column/column_factories.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/column/column_factories.pxd @@ -12,7 +12,7 @@ from pylibcudf.libcudf.types cimport ( type_id, ) -from pylibcudf.libcudf.utilities.device_buffer cimport byte, device_buffer +from pylibcudf.libcudf.utilities.device_buffer cimport device_buffer from cuda.bindings.cyruntime cimport cudaStream_t from rmm.librmm.memory_resource cimport device_async_resource_ref @@ -29,7 +29,7 @@ cdef extern from "cudf/column/column_factories.hpp" namespace "cudf" nogil: cdef unique_ptr[column] make_numeric_column( data_type type, size_type size, - device_buffer[byte] mask, + device_buffer[bitmask_type] mask, size_type null_count, cudaStream_t stream, device_async_resource_ref mr @@ -45,7 +45,7 @@ cdef extern from "cudf/column/column_factories.hpp" namespace "cudf" nogil: cdef unique_ptr[column] make_fixed_point_column( data_type type, size_type size, - device_buffer[byte] mask, + device_buffer[bitmask_type] mask, size_type null_count, cudaStream_t stream, device_async_resource_ref mr) except +libcudf_exception_handler @@ -60,7 +60,7 @@ cdef extern from "cudf/column/column_factories.hpp" namespace "cudf" nogil: cdef unique_ptr[column] make_timestamp_column( data_type type, size_type size, - device_buffer[byte] mask, + device_buffer[bitmask_type] mask, size_type null_count, cudaStream_t stream, device_async_resource_ref mr) except +libcudf_exception_handler @@ -75,7 +75,7 @@ cdef extern from "cudf/column/column_factories.hpp" namespace "cudf" nogil: cdef unique_ptr[column] make_duration_column( data_type type, size_type size, - device_buffer[byte] mask, + device_buffer[bitmask_type] mask, size_type null_count, cudaStream_t stream, device_async_resource_ref mr) except +libcudf_exception_handler @@ -90,7 +90,7 @@ cdef extern from "cudf/column/column_factories.hpp" namespace "cudf" nogil: cdef unique_ptr[column] make_fixed_width_column( data_type type, size_type size, - device_buffer[byte] mask, + device_buffer[bitmask_type] mask, size_type null_count, cudaStream_t stream, device_async_resource_ref mr) except +libcudf_exception_handler diff --git a/python/pylibcudf/pylibcudf/libcudf/null_mask.pxd b/python/pylibcudf/pylibcudf/libcudf/null_mask.pxd index 92870cfda5ed..730642cce21e 100644 --- a/python/pylibcudf/pylibcudf/libcudf/null_mask.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/null_mask.pxd @@ -7,20 +7,20 @@ from pylibcudf.exception_handler cimport libcudf_exception_handler from pylibcudf.libcudf.column.column_view cimport column_view from pylibcudf.libcudf.table.table_view cimport table_view from pylibcudf.libcudf.types cimport bitmask_type, mask_state, size_type -from pylibcudf.libcudf.utilities.device_buffer cimport byte, device_buffer +from pylibcudf.libcudf.utilities.device_buffer cimport device_buffer from cuda.bindings.cyruntime cimport cudaStream_t from rmm.librmm.memory_resource cimport device_async_resource_ref cdef extern from "cudf/null_mask.hpp" namespace "cudf" nogil: - cdef device_buffer[byte] copy_bitmask ( + cdef device_buffer[bitmask_type] copy_bitmask ( column_view view, cudaStream_t stream, device_async_resource_ref mr ) except +libcudf_exception_handler - cdef device_buffer[byte] copy_bitmask ( + cdef device_buffer[bitmask_type] copy_bitmask ( const bitmask_type* null_mask, size_type begin_bit, size_type end_bit, @@ -33,20 +33,20 @@ cdef extern from "cudf/null_mask.hpp" namespace "cudf" nogil: size_t padding_boundary ) except +libcudf_exception_handler - cdef device_buffer[byte] create_null_mask ( + cdef device_buffer[bitmask_type] create_null_mask ( size_type size, mask_state state, cudaStream_t stream, device_async_resource_ref mr ) except +libcudf_exception_handler - cdef pair[device_buffer[byte], size_type] bitmask_and( + cdef pair[device_buffer[bitmask_type], size_type] bitmask_and( table_view view, cudaStream_t stream, device_async_resource_ref mr ) - cdef pair[device_buffer[byte], size_type] bitmask_or( + cdef pair[device_buffer[bitmask_type], size_type] bitmask_or( table_view view, cudaStream_t stream, device_async_resource_ref mr @@ -77,7 +77,7 @@ cdef extern from * namespace "pylibcudf" nogil: cudf::column_view view, cudaStream_t stream, rmm::device_async_resource_ref mr) { - return std::make_unique>( + return std::make_unique>( cudf::copy_bitmask(view, stream, mr)); } @@ -86,7 +86,7 @@ cdef extern from * namespace "pylibcudf" nogil: cudf::size_type end_bit, cudaStream_t stream, rmm::device_async_resource_ref mr) { - return std::make_unique>( + return std::make_unique>( cudf::copy_bitmask(mask, begin_bit, end_bit, stream, mr)); } @@ -94,7 +94,7 @@ cdef extern from * namespace "pylibcudf" nogil: cudf::size_type size, cudf::mask_state state, cudaStream_t stream, rmm::device_async_resource_ref mr) { - return std::make_unique>( + return std::make_unique>( cudf::create_null_mask(size, state, stream, mr)); } @@ -104,7 +104,7 @@ cdef extern from * namespace "pylibcudf" nogil: { auto [mask, count] = cudf::bitmask_and(view, stream, mr); return std::pair{ - std::make_unique>(std::move(mask)), count}; + std::make_unique>(std::move(mask)), count}; } inline auto bitmask_or_unique_ptr( @@ -113,24 +113,24 @@ cdef extern from * namespace "pylibcudf" nogil: { auto [mask, count] = cudf::bitmask_or(view, stream, mr); return std::pair{ - std::make_unique>(std::move(mask)), count}; + std::make_unique>(std::move(mask)), count}; } } // namespace pylibcudf """ - cdef unique_ptr[device_buffer[byte]] copy_bitmask_to_unique_ptr( + cdef unique_ptr[device_buffer[bitmask_type]] copy_bitmask_to_unique_ptr( column_view view, cudaStream_t stream, device_async_resource_ref mr ) except +libcudf_exception_handler - cdef unique_ptr[device_buffer[byte]] copy_bitmask_to_unique_ptr( + cdef unique_ptr[device_buffer[bitmask_type]] copy_bitmask_to_unique_ptr( const bitmask_type* mask, size_type begin_bit, size_type end_bit, cudaStream_t stream, device_async_resource_ref mr ) except +libcudf_exception_handler - cdef unique_ptr[device_buffer[byte]] create_null_mask_unique_ptr( + cdef unique_ptr[device_buffer[bitmask_type]] create_null_mask_unique_ptr( size_type size, mask_state state, cudaStream_t stream, device_async_resource_ref mr ) except +libcudf_exception_handler - cdef pair[unique_ptr[device_buffer[byte]], size_type] bitmask_and_unique_ptr( + cdef pair[unique_ptr[device_buffer[bitmask_type]], size_type] bitmask_and_unique_ptr( table_view view, cudaStream_t stream, device_async_resource_ref mr ) except +libcudf_exception_handler - cdef pair[unique_ptr[device_buffer[byte]], size_type] bitmask_or_unique_ptr( + cdef pair[unique_ptr[device_buffer[bitmask_type]], size_type] bitmask_or_unique_ptr( table_view view, cudaStream_t stream, device_async_resource_ref mr ) except +libcudf_exception_handler diff --git a/python/pylibcudf/pylibcudf/libcudf/transform.pxd b/python/pylibcudf/pylibcudf/libcudf/transform.pxd index 287c67ce6209..8b7ec4a25242 100644 --- a/python/pylibcudf/pylibcudf/libcudf/transform.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/transform.pxd @@ -22,7 +22,7 @@ from pylibcudf.libcudf.types cimport ( udf_source_type, ) -from pylibcudf.libcudf.utilities.device_buffer cimport byte, device_buffer +from pylibcudf.libcudf.utilities.device_buffer cimport device_buffer from cuda.bindings.cyruntime cimport cudaStream_t from rmm.librmm.memory_resource cimport device_async_resource_ref @@ -49,7 +49,7 @@ cdef extern from "cudf/transform.hpp" namespace "cudf" nogil: ctypedef const transform_output const_transform_output - cdef pair[unique_ptr[device_buffer[byte]], size_type] bools_to_mask ( + cdef pair[unique_ptr[device_buffer[bitmask_type]], size_type] bools_to_mask ( const column_view& input, cudaStream_t stream, device_async_resource_ref mr diff --git a/python/pylibcudf/pylibcudf/null_mask.pyx b/python/pylibcudf/pylibcudf/null_mask.pyx index c448c885e5af..9dcf818944c2 100644 --- a/python/pylibcudf/pylibcudf/null_mask.pyx +++ b/python/pylibcudf/pylibcudf/null_mask.pyx @@ -9,7 +9,7 @@ from pylibcudf.libcudf cimport null_mask as cpp_null_mask from pylibcudf.libcudf.column.column_view cimport column_view from pylibcudf.libcudf.table.table_view cimport table_view from pylibcudf.libcudf.types cimport mask_state, size_type, bitmask_type -from pylibcudf.libcudf.utilities.device_buffer cimport byte, device_buffer +from pylibcudf.libcudf.utilities.device_buffer cimport device_buffer from rmm.pylibrmm.stream cimport Stream from rmm.pylibrmm.memory_resource cimport DeviceMemoryResource @@ -62,7 +62,7 @@ cpdef gpumemoryview copy_bitmask( A view containing ``col``'s bitmask, or an empty view if ``col`` is not nullable. """ - cdef unique_ptr[device_buffer[byte]] db + cdef unique_ptr[device_buffer[bitmask_type]] db cdef Stream _stream = _get_stream(stream) cdef cudaStream_t _cs = _stream.view().get() mr = _get_memory_resource(mr) @@ -109,7 +109,7 @@ cpdef gpumemoryview copy_bitmask_from_bitmask( f"bitmask must satisfy Span protocol (have .ptr and .size), " f"got {type(bitmask).__name__}" ) - cdef unique_ptr[device_buffer[byte]] db + cdef unique_ptr[device_buffer[bitmask_type]] db cdef Stream _stream = _get_stream(stream) cdef cudaStream_t _cs = _stream.view().get() mr = _get_memory_resource(mr) @@ -177,7 +177,7 @@ cpdef gpumemoryview create_null_mask( gpumemoryview A view of a null bitmask satisfying the desired size and state. """ - cdef unique_ptr[device_buffer[byte]] db + cdef unique_ptr[device_buffer[bitmask_type]] db cdef Stream _stream = _get_stream(stream) cdef cudaStream_t _cs = _stream.view().get() mr = _get_memory_resource(mr) @@ -214,7 +214,7 @@ cpdef tuple[gpumemoryview, int] bitmask_and( A tuple of the resulting mask and count of unset bits """ cdef Table c_table = Table(columns) - cdef pair[unique_ptr[device_buffer[byte]], size_type] c_result + cdef pair[unique_ptr[device_buffer[bitmask_type]], size_type] c_result cdef Stream _stream = _get_stream(stream) cdef cudaStream_t _cs = _stream.view().get() mr = _get_memory_resource(mr) @@ -255,7 +255,7 @@ cpdef tuple[gpumemoryview, int] bitmask_or( A tuple of the resulting mask and count of unset bits """ cdef Table c_table = Table(columns) - cdef pair[unique_ptr[device_buffer[byte]], size_type] c_result + cdef pair[unique_ptr[device_buffer[bitmask_type]], size_type] c_result cdef Stream _stream = _get_stream(stream) cdef cudaStream_t _cs = _stream.view().get() mr = _get_memory_resource(mr) diff --git a/python/pylibcudf/pylibcudf/nvtext/deduplicate.pyx b/python/pylibcudf/pylibcudf/nvtext/deduplicate.pyx index aaff403ec2be..c6844955002a 100644 --- a/python/pylibcudf/pylibcudf/nvtext/deduplicate.pyx +++ b/python/pylibcudf/pylibcudf/nvtext/deduplicate.pyx @@ -1,6 +1,7 @@ # SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 +from pylibcudf.libcudf.types cimport bitmask_type from cython.operator import dereference from libcpp.memory cimport unique_ptr, make_unique @@ -16,7 +17,7 @@ from pylibcudf.libcudf.nvtext.deduplicate cimport ( resolve_duplicates_pair as cpp_resolve_duplicates_pair, ) from pylibcudf.libcudf.types cimport mask_state, size_type -from pylibcudf.libcudf.utilities.device_buffer cimport byte, device_buffer +from pylibcudf.libcudf.utilities.device_buffer cimport device_buffer from pylibcudf.utils cimport _get_stream, _get_memory_resource from typing import TYPE_CHECKING @@ -36,7 +37,7 @@ cdef Column _column_from_suffix_array( cpp_suffix_array_type suffix_array, Stream stream, DeviceMemoryResource mr ): # helper to convert a suffix array to a Column - cdef unique_ptr[device_buffer[byte]] mask = ( + cdef unique_ptr[device_buffer[bitmask_type]] mask = ( cpp_null_mask.create_null_mask_unique_ptr( 0, mask_state.UNALLOCATED, diff --git a/python/pylibcudf/pylibcudf/transform.pyx b/python/pylibcudf/pylibcudf/transform.pyx index 51e66a756094..48f8c1d4ea24 100644 --- a/python/pylibcudf/pylibcudf/transform.pyx +++ b/python/pylibcudf/pylibcudf/transform.pyx @@ -28,7 +28,7 @@ from rmm.pylibrmm.memory_resource cimport DeviceMemoryResource from .column cimport Column from .expressions cimport Expression from .gpumemoryview cimport gpumemoryview, _from_cuda_device_buffer -from pylibcudf.libcudf.utilities.device_buffer cimport byte, device_buffer +from pylibcudf.libcudf.utilities.device_buffer cimport device_buffer from .types cimport DataType, null_aware, output_nullability from .utils cimport _get_stream, _get_memory_resource from typing import TYPE_CHECKING @@ -187,7 +187,7 @@ cpdef tuple[gpumemoryview, int] bools_to_mask( tuple[gpumemoryview, int] Two-tuple of a gpumemoryview wrapping the bitmask and the null count. """ - cdef pair[unique_ptr[device_buffer[byte]], size_type] c_result + cdef pair[unique_ptr[device_buffer[bitmask_type]], size_type] c_result cdef Stream _stream = _get_stream(stream) cdef cudaStream_t _cs = _stream.view().get()