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

cpp/benchmarks/bitmask/set_null_mask.cpp

@@ -39,11 +39,14 @@ auto generate_test_data(cudf::size_type num_masks,
 
   auto valids = thrust::host_vector<bool>(num_masks, true);
 
-  std::vector<rmm::device_buffer> masks(num_masks);
-  std::vector<cudf::bitmask_type*> masks_ptr(num_masks);
+  std::vector<cuda::device_buffer<std::byte>> 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) {
-    masks[i]     = cudf::create_null_mask(mask_size, cudf::mask_state::UNINITIALIZED);
-    masks_ptr[i] = static_cast<cudf::bitmask_type*>(masks[i].data());
+    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()));
   }
 
   return std::make_tuple(std::move(begin_bits),
@@ -57,19 +60,20 @@ 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"));
-  rmm::device_buffer mask = cudf::create_null_mask(mask_size, cudf::mask_state::UNINITIALIZED);
+  auto const mask_size = static_cast<cudf::size_type>(state.get_int64("mask_size"));
+  cuda::device_buffer<std::byte> 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(static_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(reinterpret_cast<cudf::bitmask_type*>(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

@@ -516,7 +516,9 @@ 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<rmm::device_buffer>(), 0};
+      : std::pair{std::make_unique<cuda::device_buffer<std::byte>>(
+                    cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)),
+                  0};
 
   return cudf::make_strings_column(num_rows,
                                    std::move(offsets),
@@ -621,7 +623,9 @@ 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<rmm::device_buffer>(), 0};
+      : std::pair{std::make_unique<cuda::device_buffer<std::byte>>(
+                    cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)),
+                  0};
 
   return std::make_unique<cudf::column>(
     dtype, num_rows, data.release(), std::move(*result_bitmask.release()), null_count);
@@ -700,7 +704,9 @@ 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<rmm::device_buffer>(), 0};
+        return std::pair{std::make_unique<cuda::device_buffer<std::byte>>(
+                           cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)),
+                         0};
       }();
 
       // Adopt remaining children as evenly as possible
@@ -785,15 +791,19 @@ std::unique_ptr<cudf::column> create_random_column<cudf::list_view>(data_profile
     thrust::device_pointer_cast(offsets.end())[-1] =
       current_child_column->size();  // Always include all elements
 
-    auto offsets_column = std::make_unique<cudf::column>(cudf::data_type{cudf::type_id::INT32},
-                                                         current_num_rows + 1,
-                                                         offsets.release(),
-                                                         rmm::device_buffer{},
-                                                         0);
+    auto offsets_column =
+      std::make_unique<cudf::column>(cudf::data_type{cudf::type_id::INT32},
+                                     current_num_rows + 1,
+                                     offsets.release(),
+                                     cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED),
+                                     0);
 
-    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<rmm::device_buffer>(), 0};
+    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>>(
+                      cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)),
+                    0};
 
     list_column = cudf::make_lists_column(current_num_rows,
                                           std::move(offsets_column),
@@ -869,8 +879,12 @@ std::unique_ptr<cudf::column> create_distinct_rows_column<cudf::list_view>(
   auto child_column      = cudf::sequence(num_rows, *zero);
   for (int lvl = dist_params.max_depth; lvl > 0; --lvl) {
     auto offsets_column = cudf::sequence(num_rows + 1, *zero);
-    auto list_column    = cudf::make_lists_column(
-      num_rows, std::move(offsets_column), std::move(child_column), 0, rmm::device_buffer{});
+    auto list_column =
+      cudf::make_lists_column(num_rows,
+                              std::move(offsets_column),
+                              std::move(child_column),
+                              0,
+                              cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED));
     if (auto const cv = list_column->view();
         cudf::has_nonempty_nulls(cv, cudf::get_default_stream())) {
       list_column = cudf::purge_nonempty_nulls(
@@ -900,8 +914,8 @@ std::unique_ptr<cudf::column> create_distinct_rows_column<cudf::struct_view>(
   children.push_back(cudf::sequence(num_rows, *cudf::make_fixed_width_scalar<int32_t>(0)));
   for (int lvl = dist_params.max_depth; lvl > 1; --lvl) {
     std::vector<std::unique_ptr<cudf::column>> parents;
-    parents.push_back(
-      cudf::create_structs_hierarchy(num_rows, std::move(children), 0, rmm::device_buffer{}));
+    parents.push_back(cudf::create_structs_hierarchy(
+      num_rows, std::move(children), 0, cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)));
     std::swap(parents, children);
   }
   auto const null_count = col->null_count();
@@ -1065,10 +1079,12 @@ std::unique_ptr<cudf::column> create_string_column(cudf::size_type num_rows,
   return std::move(table->release().front());
 }
 
-std::pair<rmm::device_buffer, cudf::size_type> create_random_null_mask(
+std::pair<cuda::device_buffer<std::byte>, cudf::size_type> create_random_null_mask(
   cudf::size_type size, std::optional<double> null_probability, unsigned seed)
 {
-  if (not null_probability.has_value()) { return {rmm::device_buffer{}, 0}; }
+  if (not null_probability.has_value()) {
+    return {cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), 0};
+  }
   CUDF_EXPECTS(*null_probability >= 0.0 and *null_probability <= 1.0,
                "Null probability must be within the range [0.0, 1.0]");
   if (*null_probability == 0.0f) {

cpp/benchmarks/common/generate_input.hpp

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

cpp/benchmarks/common/ndsh_data_generator/random_column_generator.cu

@@ -99,8 +99,11 @@ std::unique_ptr<cudf::column> generate_random_string_column(cudf::size_type lowe
                      num_rows,
                      random_string_generator(chars.data()));
 
-  return cudf::make_strings_column(
-    num_rows, std::move(offsets_column), chars.release(), 0, rmm::device_buffer{});
+  return cudf::make_strings_column(num_rows,
+                                   std::move(offsets_column),
+                                   chars.release(),
+                                   0,
+                                   cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED));
 }
 
 template <typename T>

cpp/benchmarks/copying/copy_if_else.cpp

@@ -24,9 +24,9 @@ static void bench_copy_if_else(nvbench::state& state, nvbench::type_list<DataTyp
   auto const input = create_random_table({input_type, input_type, bool_type}, row_count{num_rows});
 
   if (!nulls) {
-    input->get_column(0).set_null_mask(rmm::device_buffer{}, 0);
-    input->get_column(1).set_null_mask(rmm::device_buffer{}, 0);
-    input->get_column(2).set_null_mask(rmm::device_buffer{}, 0);
+    input->get_column(0).set_null_mask(cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), 0);
+    input->get_column(1).set_null_mask(cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), 0);
+    input->get_column(2).set_null_mask(cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), 0);
   }
 
   cudf::column_view lhs(input->view().column(0));

cpp/benchmarks/groupby/group_struct_values.cpp

@@ -41,7 +41,10 @@ static void bench_groupby_min_struct(nvbench::state& state)
 
   auto const keys_view = data_cols.front()->view();
   auto const values =
-    cudf::make_structs_column(keys_view.size(), std::move(data_cols), 0, rmm::device_buffer());
+    cudf::make_structs_column(keys_view.size(),
+                              std::move(data_cols),
+                              0,
+                              cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED));
 
   auto gb_obj   = cudf::groupby::groupby(cudf::table_view({keys_view}));
   auto requests = std::vector<cudf::groupby::aggregation_request>();
@@ -71,7 +74,10 @@ static void bench_groupby_min_struct_scan(nvbench::state& state)
 
   auto const keys_view = data_cols.front()->view();
   auto const values =
-    cudf::make_structs_column(keys_view.size(), std::move(data_cols), 0, rmm::device_buffer());
+    cudf::make_structs_column(keys_view.size(),
+                              std::move(data_cols),
+                              0,
+                              cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED));
 
   auto requests = std::vector<cudf::groupby::scan_request>();
   requests.emplace_back(cudf::groupby::scan_request());

cpp/benchmarks/io/parquet/experimental/variant/extract.cpp

@@ -288,21 +288,35 @@ std::unique_ptr<cudf::column> build_variant_column(std::span<std::span<uint8_t c
       rmm::device_buffer{offsets.data(), offsets.size() * sizeof(int32_t), stream, mr};
     auto d_data = rmm::device_buffer{flat.data(), flat.size() * sizeof(uint8_t), stream, mr};
 
-    auto off_col = std::make_unique<cudf::column>(
-      cudf::data_type{cudf::type_id::INT32}, n + 1, std::move(d_offsets), rmm::device_buffer{}, 0);
-    auto data_col = std::make_unique<cudf::column>(cudf::data_type{cudf::type_id::UINT8},
-                                                   static_cast<cudf::size_type>(flat.size()),
-                                                   std::move(d_data),
-                                                   rmm::device_buffer{},
-                                                   0);
-
-    return cudf::make_lists_column(n, std::move(off_col), std::move(data_col), 0, {});
+    auto off_col =
+      std::make_unique<cudf::column>(cudf::data_type{cudf::type_id::INT32},
+                                     n + 1,
+                                     std::move(d_offsets),
+                                     cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED),
+                                     0);
+    auto data_col =
+      std::make_unique<cudf::column>(cudf::data_type{cudf::type_id::UINT8},
+                                     static_cast<cudf::size_type>(flat.size()),
+                                     std::move(d_data),
+                                     cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED),
+                                     0);
+
+    return cudf::make_lists_column(n,
+                                   std::move(off_col),
+                                   std::move(data_col),
+                                   0,
+                                   cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED));
   };
 
   std::vector<std::unique_ptr<cudf::column>> children;
   children.emplace_back(build_list_col(meta_rows));
   children.emplace_back(build_list_col(val_rows));
-  return cudf::make_structs_column(n, std::move(children), 0, {}, stream, mr);
+  return cudf::make_structs_column(n,
+                                   std::move(children),
+                                   0,
+                                   cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED),
+                                   stream,
+                                   mr);
 }
 
 // Keys for the shared metadata dictionary: a=0, b=1, ... plus "z" for miss rows.

cpp/benchmarks/json/json.cu

@@ -162,7 +162,11 @@ auto build_json_string_column(int desired_bytes, int num_rows)
     desired_bytes, num_rows, {*d_books, *d_bicycles}, *d_book_pct, *d_misc_order, *d_store_order};
   auto [offsets, chars] = cudf::strings::detail::make_strings_children(
     jb, num_rows, cudf::get_default_stream(), cudf::get_current_device_resource_ref());
-  return cudf::make_strings_column(num_rows, std::move(offsets), chars.release(), 0, {});
+  return cudf::make_strings_column(num_rows,
+                                   std::move(offsets),
+                                   chars.release(),
+                                   0,
+                                   cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED));
 }
 
 static std::string queries[] = {"$",

cpp/benchmarks/quantiles/tdigest.cpp

@@ -50,8 +50,11 @@ void bm_tdigest_merge(nvbench::state& state)
       return i * tdigest_size;
     }));
   cudf::test::fixed_width_column_wrapper<int> offsets(offset_iter, offset_iter + num_tdigests + 1);
-  auto list_col =
-    cudf::make_lists_column(num_tdigests, offsets.release(), inner_struct.release(), 0, {});
+  auto list_col = cudf::make_lists_column(num_tdigests,
+                                          offsets.release(),
+                                          inner_struct.release(),
+                                          0,
+                                          cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED));
 
   // min and max columns
   auto min_iter = cuda::make_constant_iterator(base_value);

cpp/benchmarks/replace/clamp.cpp

@@ -25,7 +25,8 @@ void bench_clamp(nvbench::state& state, nvbench::type_list<ClampType>)
 
   auto const dtype = cudf::type_to_id<ClampType>();
   auto const input = create_random_column(dtype, row_count{n_rows});
-  if (!include_nulls) input->set_null_mask(rmm::device_buffer{}, 0);
+  if (!include_nulls)
+    input->set_null_mask(cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), 0);
 
   auto [low_scalar, high_scalar] = cudf::minmax(*input);
 

cpp/benchmarks/replace/nans.cpp

@@ -25,7 +25,8 @@ void bench_replace_nans(nvbench::state& state, nvbench::type_list<FloatingType>)
 
   auto const dtype = cudf::type_to_id<FloatingType>();
   auto const input = create_random_column(dtype, row_count{n_rows});
-  if (!include_nulls) input->set_null_mask(rmm::device_buffer{}, 0);
+  if (!include_nulls)
+    input->set_null_mask(cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED), 0);
 
   auto zero = cudf::make_fixed_width_scalar<FloatingType>(0);
 

cpp/benchmarks/rolling/rolling_sum.cpp

@@ -73,7 +73,7 @@ void bench_row_variable_rolling_sum(nvbench::state& state, nvbench::type_list<Ty
     return std::make_unique<cudf::column>(cudf::data_type(cudf::type_to_id<cudf::size_type>()),
                                           num_rows,
                                           std::move(buf),
-                                          rmm::device_buffer{},
+                                          cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED),
                                           0);
   }();
 
@@ -89,7 +89,7 @@ void bench_row_variable_rolling_sum(nvbench::state& state, nvbench::type_list<Ty
     return std::make_unique<cudf::column>(cudf::data_type(cudf::type_to_id<cudf::size_type>()),
                                           num_rows,
                                           std::move(buf),
-                                          rmm::device_buffer{},
+                                          cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED),
                                           0);
   }();