Use cudf::bitmask_type for null masks

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<cuda::device_buffer<std::byte>> masks;
+  std::vector<cuda::device_buffer<cudf::bitmask_type>> masks;
   std::vector<cudf::bitmask_type*> 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<cudf::bitmask_type*>(mask_pair.first.data()));
+    mask_pointers.push_back(mask_pair.first.data());
     return std::move(mask_pair.first);
   });
 

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<cudf::size_type>(state.get_int64("num_rows"));
   auto const mask     = cudf::create_null_mask(num_rows, cudf::mask_state::ALL_VALID);
-  auto const bitmask  = reinterpret_cast<cudf::bitmask_type const*>(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<nvbench::int8_t>(cudf::bitmask_allocation_size_bytes(num_rows));

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<bool>(num_masks, true);
 
-  std::vector<cuda::device_buffer<std::byte>> masks;
+  std::vector<cuda::device_buffer<cudf::bitmask_type>> masks;
   masks.reserve(num_masks);
   std::vector<cudf::bitmask_type*> 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<cudf::bitmask_type*>(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<cudf::size_type>(state.get_int64("mask_size"));
-  cuda::device_buffer<std::byte> mask =
+  cuda::device_buffer<cudf::bitmask_type> 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<cudf::bitmask_type*>(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");

cpp/benchmarks/common/generate_input.cu

@@ -527,7 +527,7 @@ std::unique_ptr<cudf::column> 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<bool const>(null_mask), stream)
-      : std::pair{std::make_unique<cuda::device_buffer<std::byte>>(
+      : std::pair{std::make_unique<cuda::device_buffer<cudf::bitmask_type>>(
                     cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)),
                   0};
 
@@ -640,7 +640,7 @@ std::unique_ptr<cudf::column> 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<bool const>(null_mask))
-      : std::pair{std::make_unique<cuda::device_buffer<std::byte>>(
+      : std::pair{std::make_unique<cuda::device_buffer<cudf::bitmask_type>>(
                     cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)),
                   0};
 
@@ -720,7 +720,7 @@ std::unique_ptr<cudf::column> create_random_column<cudf::struct_view>(data_profi
           return cudf::bools_to_mask(cudf::device_span<bool const>(valids),
                                      cudf::get_default_stream());
         }
-        return std::pair{std::make_unique<cuda::device_buffer<std::byte>>(
+        return std::pair{std::make_unique<cuda::device_buffer<cudf::bitmask_type>>(
                            cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)),
                          0};
       }();
@@ -817,7 +817,7 @@ std::unique_ptr<cudf::column> create_random_column<cudf::list_view>(data_profile
     auto [null_mask, null_count] =
       profile.get_null_probability().has_value()
         ? cudf::bools_to_mask(cudf::device_span<bool const>(valids))
-        : std::pair{std::make_unique<cuda::device_buffer<std::byte>>(
+        : std::pair{std::make_unique<cuda::device_buffer<cudf::bitmask_type>>(
                       cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)),
                     0};
 
@@ -1099,7 +1099,7 @@ std::unique_ptr<cudf::column> create_string_column(cudf::size_type num_rows,
   return std::move(table->release().front());
 }
 
-std::pair<cuda::device_buffer<std::byte>, cudf::size_type> create_random_null_mask(
+std::pair<cuda::device_buffer<cudf::bitmask_type>, cudf::size_type> create_random_null_mask(
   cudf::size_type size, std::optional<double> null_probability, unsigned seed)
 {
   if (not null_probability.has_value()) {

cpp/benchmarks/common/generate_input.hpp

@@ -674,5 +674,5 @@ std::vector<cudf::type_id> mix_dtypes(std::pair<cudf::type_id, cudf::type_id> 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<cuda::device_buffer<std::byte>, cudf::size_type> create_random_null_mask(
+std::pair<cuda::device_buffer<cudf::bitmask_type>, cudf::size_type> create_random_null_mask(
   cudf::size_type size, std::optional<double> null_probability = std::nullopt, unsigned seed = 1);

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<cuda::device_buffer<std::byte>, size_type> scalar_col_valid_mask_and(
+std::pair<cuda::device_buffer<cudf::bitmask_type>, size_type> scalar_col_valid_mask_and(
   column_view const& col,
   scalar const& s,
   cuda::stream_ref stream           = cudf::get_default_stream(),

cpp/include/cudf/column/column.hpp

@@ -74,7 +74,7 @@ class column {
    */
   template <typename T, CUDF_ENABLE_IF(cudf::is_numeric<T>() or cudf::is_chrono<T>())>
   column(rmm::device_uvector<T>&& other,
-         cuda::device_buffer<std::byte>&& null_mask,
+         cuda::device_buffer<cudf::bitmask_type>&& null_mask,
          size_type null_count)
     : _type{cudf::data_type{cudf::type_to_id<T>()}},
       _size{[&]() {
@@ -105,7 +105,7 @@ class column {
    * @param null_count Optional, the count of null elements.
    * @param children Optional, vector of child columns
    */
-  template <typename B1, typename B2 = cuda::device_buffer<std::byte>>
+  template <typename B1, typename B2 = cuda::device_buffer<cudf::bitmask_type>>
   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<std::byte>&& new_null_mask, size_type new_null_count);
+  void set_null_mask(cuda::device_buffer<cudf::bitmask_type>&& 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<std::byte> const& new_null_mask,
+  void set_null_mask(cuda::device_buffer<cudf::bitmask_type> 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<rmm::device_buffer> data;                   ///< data device memory buffer
-    std::unique_ptr<cuda::device_buffer<std::byte>> null_mask;  ///< null mask device memory buffer
-    std::vector<std::unique_ptr<column>> children;              ///< child columns
+    std::unique_ptr<rmm::device_buffer> data;  ///< data device memory buffer
+    std::unique_ptr<cuda::device_buffer<cudf::bitmask_type>>
+      null_mask;                                    ///< null mask device memory buffer
+    std::vector<std::unique_ptr<column>> 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<std::byte> _null_mask = cudf::create_null_mask(
+  cuda::device_buffer<cudf::bitmask_type> _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

cpp/include/cudf/column/column_factories.hpp

@@ -442,7 +442,7 @@ std::unique_ptr<column> make_strings_column(size_type num_strings,
                                             std::unique_ptr<column> offsets_column,
                                             rmm::device_buffer&& chars_buffer,
                                             size_type null_count,
-                                            cuda::device_buffer<std::byte>&& null_mask);
+                                            cuda::device_buffer<cudf::bitmask_type>&& null_mask);
 
 /**
  * @brief Construct a LIST type column given offsets column, child column, null mask and null
@@ -497,11 +497,12 @@ std::unique_ptr<column> 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<cudf::column> make_lists_column(size_type num_rows,
-                                                std::unique_ptr<column> offsets_column,
-                                                std::unique_ptr<column> child_column,
-                                                size_type null_count,
-                                                cuda::device_buffer<std::byte>&& null_mask);
+std::unique_ptr<cudf::column> make_lists_column(
+  size_type num_rows,
+  std::unique_ptr<column> offsets_column,
+  std::unique_ptr<column> child_column,
+  size_type null_count,
+  cuda::device_buffer<cudf::bitmask_type>&& null_mask);
 
 /**
  * @brief Create an empty LIST column
@@ -540,7 +541,7 @@ std::unique_ptr<cudf::column> make_structs_column(
   size_type num_rows,
   std::vector<std::unique_ptr<column>>&& child_columns,
   size_type null_count,
-  cuda::device_buffer<std::byte>&& null_mask,
+  cuda::device_buffer<cudf::bitmask_type>&& 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<cudf::column> create_structs_hierarchy(
   size_type num_rows,
   std::vector<std::unique_ptr<column>>&& child_columns,
   size_type null_count,
-  cuda::device_buffer<std::byte>&& null_mask,
+  cuda::device_buffer<cudf::bitmask_type>&& null_mask,
   cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 

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<std::byte>`
+ * `views` into a `cuda::device_buffer<cudf::bitmask_type>`
  *
  * 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<std::byte> concatenate_masks(
+cuda::device_buffer<cudf::bitmask_type> concatenate_masks(
   std::span<column_view const> views,
   cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());

cpp/include/cudf/detail/concatenate_masks.hpp

@@ -56,9 +56,9 @@ size_type concatenate_masks(host_span<column_view const> views,
  *
  * @param stream CUDA stream used for device memory operations and kernel launches.
  */
-cuda::device_buffer<std::byte> concatenate_masks(std::span<column_view const> views,
-                                                 cuda::stream_ref stream,
-                                                 rmm::device_async_resource_ref mr);
+cuda::device_buffer<cudf::bitmask_type> concatenate_masks(std::span<column_view const> views,
+                                                          cuda::stream_ref stream,
+                                                          rmm::device_async_resource_ref mr);
 
 }  // namespace detail
 }  // namespace cudf

cpp/include/cudf/detail/null_mask.cuh

@@ -256,7 +256,7 @@ rmm::device_uvector<size_type> inplace_segmented_bitmask_binop(
  * @param stream CUDA stream used for device memory operations and kernel launches
  */
 template <typename Binop>
-std::pair<cuda::device_buffer<std::byte>, size_type> bitmask_binop(
+std::pair<cuda::device_buffer<cudf::bitmask_type>, size_type> bitmask_binop(
   Binop op,
   host_span<bitmask_type const* const> masks,
   host_span<size_type const> masks_begin_bits,
@@ -267,21 +267,21 @@ std::pair<cuda::device_buffer<std::byte>, 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<bitmask_type>(reinterpret_cast<bitmask_type*>(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<bitmask_type>(
+                                             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 <typename Binop>
-std::pair<std::vector<std::unique_ptr<cuda::device_buffer<std::byte>>>, std::vector<size_type>>
+std::pair<std::vector<std::unique_ptr<cuda::device_buffer<cudf::bitmask_type>>>,
+          std::vector<size_type>>
 segmented_bitmask_binop(Binop op,
                         host_span<bitmask_type const* const> masks,
                         host_span<size_type const> 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<std::unique_ptr<cuda::device_buffer<std::byte>>> h_destination_masks;
+  std::vector<std::unique_ptr<cuda::device_buffer<cudf::bitmask_type>>> h_destination_masks;
   h_destination_masks.reserve(num_segments);
   std::vector<bitmask_type*> 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<cuda::device_buffer<std::byte>>(
+    h_destination_masks.push_back(std::make_unique<cuda::device_buffer<cudf::bitmask_type>>(
       cudf::create_null_mask(num_bytes * CHAR_BIT, cudf::mask_state::UNINITIALIZED, stream, mr)));
-    h_destination_masks_ptrs.push_back(
-      reinterpret_cast<bitmask_type*>(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<size_type> segmented_null_count(bitmask_type const* bitmask,
  * @return A pair containing the reduced null mask and number of nulls.
  */
 template <typename OffsetIterator>
-std::pair<cuda::device_buffer<std::byte>, size_type> segmented_null_mask_reduction(
+std::pair<cuda::device_buffer<cudf::bitmask_type>, size_type> segmented_null_mask_reduction(
   bitmask_type const* bitmask,
   OffsetIterator first_bit_indices_begin,
   OffsetIterator first_bit_indices_end,