Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Refactor bit counting APIs, introduce valid/null count functions, and split host/device side code for segmented counts. #9588

Merged
merged 57 commits into from
Dec 10, 2021
Merged
Show file tree
Hide file tree
Changes from 51 commits
Commits
Show all changes
57 commits
Select commit Hold shift + click to select a range
99ab7cc
Separate device function for segmented_count_set_bits.
bdice Oct 28, 2021
1fa8421
Simplify to_word_index_functor.
bdice Oct 28, 2021
3a7f6c3
Intermediate work.
bdice Nov 2, 2021
2f82f0d
Fix host implementation of segmented_count_set_bits.
bdice Nov 3, 2021
7c2adf5
Remove vector utility function - it doesn't make sense in this scope.
bdice Nov 3, 2021
aa184a6
Use indexing.
bdice Nov 3, 2021
cafa1e9
Move variable declaration.
bdice Nov 3, 2021
0fcf354
Move single-use functors and kernels from .cuh to .cu.
bdice Nov 3, 2021
652ebf2
Minor refactoring for constness, better comments, clarity.
bdice Nov 3, 2021
c072278
Clarify segment behavior.
bdice Nov 3, 2021
3961fb3
Add comments.
bdice Nov 3, 2021
5caeddf
Add template for std::vector.
bdice Nov 4, 2021
3c343b3
Unify detail implementations of segmented_count_set_bits and segmente…
bdice Nov 4, 2021
96d837c
Merge remote-tracking branch 'upstream/branch-21.12' into segmented-n…
bdice Nov 9, 2021
5201800
Merge remote-tracking branch 'upstream/branch-21.12' into segmented-n…
bdice Nov 11, 2021
9076630
Merge remote-tracking branch 'upstream/branch-22.02' into segmented-n…
bdice Nov 11, 2021
55f6817
Temp.
bdice Nov 16, 2021
40943fd
Refactorings.
bdice Nov 17, 2021
5a22bad
Merge remote-tracking branch 'upstream/branch-22.02' into unify-segme…
bdice Nov 17, 2021
d88b679
Minor refactoring.
bdice Nov 17, 2021
9af7e91
Switch some functions to use offset iterators.
bdice Nov 17, 2021
a5c3268
Rename index_alternator.
bdice Nov 17, 2021
fc5ab33
Revert null count update.
bdice Nov 17, 2021
a131acf
Swap order of code.
bdice Nov 17, 2021
37e2488
Minor fixes to docs, includes.
bdice Nov 17, 2021
f5930b6
Fix bug, minor refactors.
bdice Nov 17, 2021
91d0b11
Revert change in docs.
bdice Nov 17, 2021
a40dbf9
Merge branch 'unify-segmented_count_unset_bits' into segmented-null-f…
bdice Nov 17, 2021
f626d8e
Merge remote-tracking branch 'upstream/branch-22.02' into segmented-n…
bdice Nov 17, 2021
1daac06
Update variables/docs.
bdice Nov 17, 2021
fa2efc1
Use transform iterator directly.
bdice Nov 17, 2021
2f76450
Doc updates.
bdice Nov 17, 2021
b54c89c
Refactor detail APIs to expose host-side segmented functions named se…
bdice Nov 19, 2021
56d42bf
Merge remote-tracking branch 'upstream/branch-22.02' into segmented-n…
bdice Nov 23, 2021
f568bff
Refactor transform iterator.
bdice Nov 23, 2021
1f88efd
Change to begin/end style, use 'inclusive' as parameter name.
bdice Nov 24, 2021
af4698a
Add doxygen brief to count_bits_policy enum.
bdice Nov 24, 2021
65e3745
Use transform.
bdice Nov 29, 2021
cc1abc8
Merge remote-tracking branch 'upstream/branch-22.02' into segmented-n…
bdice Nov 30, 2021
8d3229f
Use make_std_vector_sync to copy data back to the host.
bdice Nov 30, 2021
33db353
Fix comments / docstrings.
bdice Dec 1, 2021
b2b50a2
Intermediate work.
bdice Dec 3, 2021
6f380c6
Intermediate work.
bdice Dec 8, 2021
78a3cfc
Remove public bit counting APIs, refactor behavior and docs, add deta…
bdice Dec 8, 2021
91a827a
Intermediate work.
bdice Dec 9, 2021
66e000e
Fix host_span functions.
bdice Dec 9, 2021
7a1ae12
Merge remote-tracking branch 'upstream/branch-22.02' into segmented-n…
bdice Dec 9, 2021
3884148
Fix constructor.
bdice Dec 9, 2021
2433275
Fix tests for bitmask==nullptr.
bdice Dec 9, 2021
1e7d35d
Use thrust::get.
bdice Dec 1, 2021
bbd045f
Merge remote-tracking branch 'upstream/branch-22.02' into segmented-n…
bdice Dec 9, 2021
b4384d3
Fix docs.
bdice Dec 9, 2021
254a6cc
Add tests for valid/null counts.
bdice Dec 9, 2021
837ebe2
Use ElementsAreArray instead of ContainerEq (order matters).
bdice Dec 9, 2021
e9d9c0e
Add missing test.
bdice Dec 9, 2021
1ee405b
Fix clang-format.
bdice Dec 9, 2021
3f689bc
Reorder includes.
bdice Dec 10, 2021
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
406 changes: 230 additions & 176 deletions cpp/include/cudf/detail/null_mask.cuh

Large diffs are not rendered by default.

131 changes: 125 additions & 6 deletions cpp/include/cudf/detail/null_mask.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -48,43 +48,162 @@ void set_null_mask(bitmask_type* bitmask,
rmm::cuda_stream_view stream = rmm::cuda_stream_default);

/**
* @copydoc cudf::count_set_bits
* @brief Given a bitmask, counts the number of set (1) bits in the range
* `[start, stop)`.
*
* @param[in] stream CUDA stream used for device memory operations and kernel launches.
* @throws cudf::logic_error if `bitmask == nullptr`
* @throws cudf::logic_error if `start > stop`
* @throws cudf::logic_error if `start < 0`
*
* @param bitmask Bitmask residing in device memory whose bits will be counted.
* @param start Index of the first bit to count (inclusive).
* @param stop Index of the last bit to count (exclusive).
* @param stream CUDA stream used for device memory operations and kernel launches.
* @return The number of non-zero bits in the specified range.
*/
cudf::size_type count_set_bits(bitmask_type const* bitmask,
size_type start,
size_type stop,
rmm::cuda_stream_view stream);

/**
* @copydoc cudf::count_unset_bits
* @brief Given a bitmask, counts the number of unset (0) bits in the range
*`[start, stop)`.
*
* @param[in] stream CUDA stream used for device memory operations and kernel launches.
* Returns `0` if `bitmask == nullptr`.
*
* @throws cudf::logic_error if `bitmask == nullptr`
* @throws cudf::logic_error if `start > stop`
* @throws cudf::logic_error if `start < 0`
*
* @param bitmask Bitmask residing in device memory whose bits will be counted.
* @param start Index of the first bit to count (inclusive).
* @param stop Index of the last bit to count (exclusive).
* @param stream CUDA stream used for device memory operations and kernel launches.
* @return The number of zero bits in the specified range.
*/
cudf::size_type count_unset_bits(bitmask_type const* bitmask,
size_type start,
size_type stop,
rmm::cuda_stream_view stream);

/**
* @copydoc cudf::segmented_count_set_bits
* @brief Given a bitmask, counts the number of set (1) bits in every range
* `[indices[2*i], indices[(2*i)+1])` (where 0 <= i < indices.size() / 2).
*
* @throws cudf::logic_error if `bitmask == nullptr`
* @throws cudf::logic_error if `indices.size() % 2 != 0`
* @throws cudf::logic_error if `indices[2*i] < 0 or indices[2*i] > indices[(2*i)+1]`
*
* @param[in] bitmask Bitmask residing in device memory whose bits will be counted.
* @param[in] indices A host_span of indices specifying ranges to count the number of set bits.
* @param[in] stream CUDA stream used for device memory operations and kernel launches.
* @return A vector storing the number of non-zero bits in the specified ranges.
*/
std::vector<size_type> segmented_count_set_bits(bitmask_type const* bitmask,
host_span<size_type const> indices,
rmm::cuda_stream_view stream);

/**
* @copydoc cudf::segmented_count_unset_bits
* @brief Given a bitmask, counts the number of unset (0) bits in every range
* `[indices[2*i], indices[(2*i)+1])` (where 0 <= i < indices.size() / 2).
*
* @throws cudf::logic_error if `bitmask == nullptr`
* @throws cudf::logic_error if `indices.size() % 2 != 0`
* @throws cudf::logic_error if `indices[2*i] < 0 or indices[2*i] > indices[(2*i)+1]`
*
* @param[in] bitmask Bitmask residing in device memory whose bits will be counted.
* @param[in] indices A host_span of indices specifying ranges to count the number of unset bits.
* @param[in] stream CUDA stream used for device memory operations and kernel launches.
* @return A vector storing the number of zero bits in the specified ranges.
*/
std::vector<size_type> segmented_count_unset_bits(bitmask_type const* bitmask,
host_span<size_type const> indices,
rmm::cuda_stream_view stream);

/**
* @brief Given a validity bitmask, counts the number of valid elements (set bits)
* in the range `[start, stop)`.
*
* If `bitmask == nullptr`, all elements are assumed to be valid and the
* function returns `stop-start`.
*
* @throws cudf::logic_error if `start > stop`
* @throws cudf::logic_error if `start < 0`
*
* @param[in] bitmask Validity bitmask residing in device memory.
* @param[in] start Index of the first bit to count (inclusive).
* @param[in] stop Index of the last bit to count (exclusive).
* @param[in] stream CUDA stream used for device memory operations and kernel launches.
* @return The number of valid elements in the specified range.
*/
cudf::size_type valid_count(bitmask_type const* bitmask,
size_type start,
size_type stop,
rmm::cuda_stream_view stream);

/**
* @brief Given a validity bitmask, counts the number of null elements (unset bits)
* in the range `[start, stop)`.
*
* If `bitmask == nullptr`, all elements are assumed to be valid and the
* function returns ``.
*
* @throws cudf::logic_error if `start > stop`
* @throws cudf::logic_error if `start < 0`
*
* @param[in] bitmask Validity bitmask residing in device memory.
* @param[in] start Index of the first bit to count (inclusive).
* @param[in] stop Index of the last bit to count (exclusive).
* @param[in] stream CUDA stream used for device memory operations and kernel launches.
* @return The number of null elements in the specified range.
*/
cudf::size_type null_count(bitmask_type const* bitmask,
size_type start,
size_type stop,
rmm::cuda_stream_view stream);

/**
* @brief Given a validity bitmask, counts the number of valid elements (set
* bits) in every range `[indices[2*i], indices[(2*i)+1])` (where 0 <= i <
* indices.size() / 2).
*
* If `bitmask == nullptr`, all elements are assumed to be valid and a vector of
* length `indices.size()` containing segment lengths is returned.
*
* @throws cudf::logic_error if `indices.size() % 2 != 0`.
* @throws cudf::logic_error if `indices[2*i] < 0 or indices[2*i] > indices[(2*i)+1]`.
*
* @param[in] bitmask Validity bitmask residing in device memory.
* @param[in] indices A host_span of indices specifying ranges to count the number of valid
* elements.
* @param[in] stream CUDA stream used for device memory operations and kernel launches.
* @return A vector storing the number of valid elements in each specified range.
*/
std::vector<size_type> segmented_valid_count(bitmask_type const* bitmask,
host_span<size_type const> indices,
rmm::cuda_stream_view stream);

/**
* @brief Given a validity bitmask, counts the number of null elements (unset
* bits) in every range `[indices[2*i], indices[(2*i)+1])` (where 0 <= i <
* indices.size() / 2).
*
* If `bitmask == nullptr`, all elements are assumed to be valid and a vector of
* length `indices.size()` containing all zeros is returned.
*
* @throws cudf::logic_error if `indices.size() % 2 != 0`
* @throws cudf::logic_error if `indices[2*i] < 0 or indices[2*i] > indices[(2*i)+1]`
*
* @param[in] bitmask Validity bitmask residing in device memory.
* @param[in] indices A host_span of indices specifying ranges to count the number of null elements.
* @param[in] stream CUDA stream used for device memory operations and kernel launches.
* @return A vector storing the number of null elements in each specified range.
*/
std::vector<size_type> segmented_null_count(bitmask_type const* bitmask,
host_span<size_type const> indices,
rmm::cuda_stream_view stream);

/**
* @copydoc cudf::copy_bitmask(bitmask_type const*, size_type, size_type,
*rmm::mr::device_memory_resource*)
Expand Down
64 changes: 0 additions & 64 deletions cpp/include/cudf/null_mask.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -100,70 +100,6 @@ rmm::device_buffer create_null_mask(
*/
void set_null_mask(bitmask_type* bitmask, size_type begin_bit, size_type end_bit, bool valid);

/**
* @brief Given a bitmask, counts the number of set (1) bits in the range
* `[start, stop)`
*
* Returns `0` if `bitmask == nullptr`.
*
* @throws cudf::logic_error if `start > stop`
* @throws cudf::logic_error if `start < 0`
*
* @param bitmask Bitmask residing in device memory whose bits will be counted
* @param start Index of the first bit to count (inclusive)
* @param stop Index of the last bit to count (exclusive)
* @return The number of non-zero bits in the specified range
*/
cudf::size_type count_set_bits(bitmask_type const* bitmask, size_type start, size_type stop);

/**
* @brief Given a bitmask, counts the number of unset (0) bits in the range
*`[start, stop)`.
*
* Returns `0` if `bitmask == nullptr`.
*
* @throws cudf::logic_error if `start > stop`
* @throws cudf::logic_error if `start < 0`
*
* @param bitmask Bitmask residing in device memory whose bits will be counted
* @param start Index of the first bit to count (inclusive)
* @param stop Index of the last bit to count (exclusive)
* @return The number of zero bits in the specified range
*/
cudf::size_type count_unset_bits(bitmask_type const* bitmask, size_type start, size_type stop);

/**
* @brief Given a bitmask, counts the number of set (1) bits in every range
* `[indices[2*i], indices[(2*i)+1])` (where 0 <= i < indices.size() / 2).
*
* Returns an empty vector if `bitmask == nullptr`.
*
* @throws cudf::logic_error if `indices.size() % 2 != 0`
* @throws cudf::logic_error if `indices[2*i] < 0 or indices[2*i] > indices[(2*i)+1]`
*
* @param[in] bitmask Bitmask residing in device memory whose bits will be counted
* @param[in] indices A host_span of indices specifying ranges to count the number of set bits
* @return A vector storing the number of non-zero bits in the specified ranges
*/
std::vector<size_type> segmented_count_set_bits(bitmask_type const* bitmask,
host_span<cudf::size_type const> indices);

/**
* @brief Given a bitmask, counts the number of unset (0) bits in every range
* `[indices[2*i], indices[(2*i)+1])` (where 0 <= i < indices.size() / 2).
*
* Returns an empty vector if `bitmask == nullptr`.
*
* @throws cudf::logic_error if `indices.size() % 2 != 0`
* @throws cudf::logic_error if `indices[2*i] < 0 or indices[2*i] > indices[(2*i)+1]`
*
* @param[in] bitmask Bitmask residing in device memory whose bits will be counted
* @param[in] indices A host_span of indices specifying ranges to count the number of unset bits
* @return A vector storing the number of zero bits in the specified ranges
*/
std::vector<size_type> segmented_count_unset_bits(bitmask_type const* bitmask,
host_span<cudf::size_type const> indices);

/**
* @brief Creates a `device_buffer` from a slice of bitmask defined by a range
* of indices `[begin_bit, end_bit)`.
Expand Down
Loading