Fix missing null_count() comparison in test framework and related failures (#7219)

Fixes #7210
Fixes #6733 

List of fixes included:

- [x] Restore `null_count()` check in `expect_columns_equal` / `expect_columns_equivalent`
- [x] Fix issue in `structs_column_view::get_sliced_child`
- [x] Fix test failures in COPYING_TEST
- [x] Fix test failures in STREAM_COMPACTION_TEST
- [x] Fix test failures in RESHAPE_TEST

Authors:
  - @nvdbaranec
  - Mark Harris (@harrism)

Approvers:
  - Mark Harris (@harrism)
  - MithunR (@mythrocks)
  - Jake Hemstad (@jrhemstad)

URL: #7219

Author: nvdbaranec

cpp/include/cudf/detail/copy_if.cuh

@@ -99,7 +99,6 @@ __launch_bounds__(block_size) __global__
 {
   T* __restrict__ output_data                   = output_view.data<T>();
   cudf::bitmask_type* __restrict__ output_valid = output_view.null_mask();
-  constexpr cudf::size_type leader_lane{0};
   static_assert(block_size <= 1024, "Maximum thread block size exceeded");
 
   int tid                      = threadIdx.x + per_thread * block_size * blockIdx.x;
@@ -109,8 +108,8 @@ __launch_bounds__(block_size) __global__
   __shared__ bool temp_valids[has_validity ? block_size + cudf::detail::warp_size : 1];
   __shared__ T temp_data[block_size];
 
-  cudf::size_type warp_valid_counts{0};
-  cudf::size_type block_sum = 0;
+  cudf::size_type warp_valid_counts{0};  // total valid sum over the `per_thread` loop below
+  cudf::size_type block_sum = 0;         // count passing filter over the `per_thread` loop below
 
   // Note that since the maximum gridDim.x on all supported GPUs is as big as
   // cudf::size_type, this loop is sufficient to cover our maximum column size
@@ -160,6 +159,8 @@ __launch_bounds__(block_size) __global__
       const int wid        = threadIdx.x / cudf::detail::warp_size;
       const int lane       = threadIdx.x % cudf::detail::warp_size;
 
+      cudf::size_type tmp_warp_valid_counts{0};
+
       if (tmp_block_sum > 0 && wid <= last_warp) {
         int valid_index = (block_offset / cudf::detail::warp_size) + wid;
 
@@ -168,9 +169,8 @@ __launch_bounds__(block_size) __global__
 
         // Note the atomicOr's below assume that output_valid has been set to
         // all zero before the kernel
-
         if (lane == 0 && valid_warp != 0) {
-          warp_valid_counts = __popc(valid_warp);
+          tmp_warp_valid_counts = __popc(valid_warp);
           if (wid > 0 && wid < last_warp)
             output_valid[valid_index] = valid_warp;
           else {
@@ -182,19 +182,22 @@ __launch_bounds__(block_size) __global__
         if ((wid == 0) && (last_warp == num_warps)) {
           uint32_t valid_warp = __ballot_sync(0xffffffff, temp_valids[block_size + threadIdx.x]);
           if (lane == 0 && valid_warp != 0) {
-            warp_valid_counts += __popc(valid_warp);
+            tmp_warp_valid_counts += __popc(valid_warp);
             atomicOr(&output_valid[valid_index + num_warps], valid_warp);
           }
         }
       }
+      warp_valid_counts += tmp_warp_valid_counts;
     }
 
     block_offset += tmp_block_sum;
     tid += block_size;
   }
   // Compute total null_count for this block and add it to global count
+  constexpr cudf::size_type leader_lane{0};
   cudf::size_type block_valid_count =
     cudf::detail::single_lane_block_sum_reduce<block_size, leader_lane>(warp_valid_counts);
+
   if (threadIdx.x == 0) {  // one thread computes and adds to null count
     atomicAdd(output_null_count, block_sum - block_valid_count);
   }

cpp/src/structs/structs_column_view.cpp

@@ -36,7 +36,9 @@ column_view structs_column_view::get_sliced_child(int index) const
                      size(),
                      child(index).head<uint8_t>(),
                      child(index).null_mask(),
-                     child(index).null_count(),
+                     // TODO: could potentially compute the actual count here, but at
+                     // the moment this interface doesn't take a stream.
+                     UNKNOWN_NULL_COUNT,
                      offset(),
                      children};
 }

cpp/tests/copying/scatter_list_tests.cu

@@ -733,7 +733,9 @@ TYPED_TEST(TypedScatterListsTest, ListsOfNullStructs)
   };
   // clang-format on
 
-  auto expected_structs = structs_column_wrapper{{expected_numerics, expected_strings}};
+  auto expected_structs =
+    structs_column_wrapper{{expected_numerics, expected_strings},
+                           make_counting_transform_iterator(0, [](auto i) { return i != 6; })};
 
   auto expected_lists = cudf::make_lists_column(
     6, offsets_column{0, 3, 5, 9, 11, 13, 15}.release(), expected_structs.release(), 0, {});
@@ -828,7 +830,9 @@ TYPED_TEST(TypedScatterListsTest, EmptyListsOfStructs)
   };
   // clang-format on
 
-  auto expected_structs = structs_column_wrapper{{expected_numerics, expected_strings}};
+  auto expected_structs =
+    structs_column_wrapper{{expected_numerics, expected_strings},
+                           make_counting_transform_iterator(0, [](auto i) { return i != 6; })};
 
   auto expected_lists = cudf::make_lists_column(
     6, offsets_column{0, 3, 5, 9, 11, 11, 13}.release(), expected_structs.release(), 0, {});
@@ -929,7 +933,9 @@ TYPED_TEST(TypedScatterListsTest, NullListsOfStructs)
   };
   // clang-format on
 
-  auto expected_structs = structs_column_wrapper{{expected_numerics, expected_strings}};
+  auto expected_structs =
+    structs_column_wrapper{{expected_numerics, expected_strings},
+                           make_counting_transform_iterator(0, [](auto i) { return i != 6; })};
 
   auto expected_lists_null_mask_begin =
     make_counting_transform_iterator(0, [](auto i) { return i != 4; });

cpp/tests/reshape/byte_cast_tests.cpp

@@ -114,7 +114,7 @@ TEST_F(ByteCastTest, int32ValuesWithNulls)
     5,
     std::move(fixed_width_column_wrapper<cudf::size_type>{0, 4, 8, 12, 16, 20}.release()),
     std::move(int32_data.release()),
-    3,
+    2,
     detail::make_null_mask(even_validity, even_validity + 5));
 
   auto const output_int32 = cudf::byte_cast(int32_col, cudf::flip_endianness::YES);

cpp/tests/stream_compaction/apply_boolean_mask_tests.cpp

@@ -249,7 +249,7 @@ TEST_F(ApplyBooleanMask, NoNullInput)
 
 TEST_F(ApplyBooleanMask, CorrectNullCount)
 {
-  cudf::size_type inputRows = 75000;
+  cudf::size_type inputRows = 471234;
 
   auto seq1       = cudf::test::make_counting_transform_iterator(0, [](auto i) { return i; });
   auto valid_seq1 = cudf::test::make_counting_transform_iterator(0, [](auto row) { return true; });

cpp/tests/utilities/column_utilities.cu

@@ -67,6 +67,8 @@ struct column_property_comparator {
 
     if (lhs.size() > 0 && check_exact_equality) { EXPECT_EQ(lhs.nullable(), rhs.nullable()); }
 
+    EXPECT_EQ(lhs.null_count(), rhs.null_count());
+
     // equivalent, but not exactly equal columns can have a different number of children if their
     // sizes are both 0. Specifically, empty string columns may or may not have children.
     if (check_exact_equality || lhs.size() > 0) {