Fix null_count of columns returned by chunked_parquet_reader

cpp/src/io/parquet/page_data.cu

@@ -1018,23 +1018,25 @@ static __device__ bool setupLocalPageInfo(page_state_s* const s,
           s->page.nesting_decode[thread_depth].start_depth;
         s->nesting_decode_cache[thread_depth].end_depth =
           s->page.nesting_decode[thread_depth].end_depth;
+        s->nesting_decode_cache[thread_depth].valid_count = 0;
+        s->nesting_decode_cache[thread_depth].value_count = 0;
+        s->nesting_decode_cache[thread_depth].null_count  = 0;
       }
       depth += blockDim.x;
     }
   }
   if (!t) {
     s->nesting_info = can_use_decode_cache ? s->nesting_decode_cache : s->page.nesting_decode;
   }
-  __syncthreads();
 
   // zero counts
   int depth = 0;
   while (depth < s->page.num_output_nesting_levels) {
     int const thread_depth = depth + t;
     if (thread_depth < s->page.num_output_nesting_levels) {
-      s->nesting_info[thread_depth].valid_count = 0;
-      s->nesting_info[thread_depth].value_count = 0;
-      s->nesting_info[thread_depth].null_count  = 0;
+      s->page.nesting_decode[thread_depth].valid_count = 0;
+      s->page.nesting_decode[thread_depth].value_count = 0;
+      s->page.nesting_decode[thread_depth].null_count  = 0;
     }
     depth += blockDim.x;
   }

cpp/src/io/utilities/column_buffer.cpp

@@ -241,13 +241,8 @@ std::unique_ptr<column> empty_like(column_buffer& buffer,
       auto child = empty_like(buffer.children[0], child_info, stream, mr);
 
       // make the final list column
-      return make_lists_column(0,
-                               std::move(offsets),
-                               std::move(child),
-                               buffer._null_count,
-                               std::move(buffer._null_mask),
-                               stream,
-                               mr);
+      return make_lists_column(
+        0, std::move(offsets), std::move(child), 0, rmm::device_buffer{0, stream, mr}, stream, mr);
     } break;
 
     case type_id::STRUCT: {
@@ -265,12 +260,8 @@ std::unique_ptr<column> empty_like(column_buffer& buffer,
                        return empty_like(col, child_info, stream, mr);
                      });
 
-      return make_structs_column(0,
-                                 std::move(output_children),
-                                 buffer._null_count,
-                                 std::move(buffer._null_mask),
-                                 stream,
-                                 mr);
+      return make_structs_column(
+        0, std::move(output_children), 0, rmm::device_buffer{0, stream, mr}, stream, mr);
     } break;
 
     default: return cudf::make_empty_column(buffer.type);

cpp/tests/io/parquet_chunked_reader_test.cpp

@@ -920,3 +920,34 @@ TEST_F(ParquetChunkedReaderTest, TestChunkedReadWithListsOfStructs)
     CUDF_TEST_EXPECT_TABLES_EQUAL(*expected_with_nulls, *result);
   }
 }
+
+TEST_F(ParquetChunkedReaderTest, TestChunkedReadNullCount)
+{
+  auto constexpr num_rows = 100'000;
+
+  auto const sequence = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return 1; });
+  auto const validity =
+    cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i % 4 != 3; });
+  cudf::test::fixed_width_column_wrapper<int32_t> col{sequence, sequence + num_rows, validity};
+  std::vector<std::unique_ptr<cudf::column>> cols;
+  cols.push_back(col.release());
+  auto const expected = std::make_unique<cudf::table>(std::move(cols));
+
+  auto const filepath        = temp_env->get_temp_filepath("chunked_reader_null_count.parquet");
+  auto const page_limit_rows = num_rows / 5;
+  auto const write_opts =
+    cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, *expected)
+      .max_page_size_rows(page_limit_rows)  // 20k rows per page
+      .build();
+  cudf::io::write_parquet(write_opts);
+
+  auto const byte_limit = page_limit_rows * sizeof(int);
+  auto const read_opts =
+    cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}).build();
+  auto reader = cudf::io::chunked_parquet_reader(byte_limit, read_opts);
+
+  do {
+    // Every fourth row is null
+    EXPECT_EQ(reader.read_chunk().tbl->get_column(0).null_count(), page_limit_rows / 4);
+  } while (reader.has_next());
+}