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.

Sign up for GitHub

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

Jump to bottom

Fix the issue of empty lists having empty offsets #10935

Merged
merged 1 commit into from
May 25, 2022
Merged

Fix the issue of empty lists having empty offsets #10935

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
45 changes: 45 additions & 0 deletions cpp/src/io/parquet/page_enc.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1613,6 +1613,51 @@ dremel_data get_dremel_data(column_view h_col,
return std::make_tuple(std::move(empties), std::move(empties_idx), empties_size);
};

// Check if there are empty lists with empty offsets in this column
bool has_empty_list_offsets = false;
{
auto curr_col = h_col;
while (is_nested(curr_col.type())) {
if (curr_col.type().id() == type_id::LIST) {
auto lcv = lists_column_view(curr_col);
if (lcv.offsets().size() == 0) {
has_empty_list_offsets = true;
break;
}
curr_col = lcv.child();
} else if (curr_col.type().id() == type_id::STRUCT) {
curr_col = curr_col.child(0);
vuule marked this conversation as resolved.
Show resolved Hide resolved
}
}
}
std::unique_ptr<column> empty_list_offset_col;
if (has_empty_list_offsets) {
empty_list_offset_col = make_fixed_width_column(data_type(type_id::INT32), 1);
cudaMemsetAsync(empty_list_offset_col->mutable_view().head(), 0, sizeof(size_type), stream);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I wonder if this could be done with a simple column_view wrapper around a pointer to
__constant__ int32_t empty_list_offset_data = 0;
This is probably not worth doing for a single value just to save the Memset call.

std::function<column_view(column_view const&)> normalize_col = [&](column_view const& col) {
auto children = [&]() -> std::vector<column_view> {
if (col.type().id() == type_id::LIST) {
auto lcol = lists_column_view(col);
auto offset_col =
lcol.offsets().head() == nullptr ? empty_list_offset_col->view() : lcol.offsets();
return {offset_col, normalize_col(lcol.child())};
} else if (col.type().id() == type_id::STRUCT) {
return {normalize_col(col.child(0))};
} else {
return {col.child_begin(), col.child_end()};
}
}();
return column_view(col.type(),
col.size(),
col.head(),
col.null_mask(),
UNKNOWN_NULL_COUNT,
col.offset(),
std::move(children));
};
h_col = normalize_col(h_col);
}

auto curr_col = h_col;
std::vector<column_view> nesting_levels;
std::vector<uint8_t> def_at_level;
Expand Down
8 changes: 4 additions & 4 deletions cpp/src/io/parquet/parquet_gpu.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
#include "io/utilities/hostdevice_vector.hpp"

#include <cudf/column/column_device_view.cuh>
#include <cudf/lists/lists_column_view.hpp>
#include <cudf/lists/lists_column_device_view.cuh>
#include <cudf/table/table_device_view.cuh>
#include <cudf/types.hpp>
#include <cudf/utilities/span.hpp>
Expand Down Expand Up @@ -305,9 +305,9 @@ inline size_type __device__ row_to_value_idx(size_type idx, column_device_view c
idx += col.offset();
col = col.child(0);
} else {
auto offset_col = col.child(lists_column_view::offsets_column_index);
idx = offset_col.element<size_type>(idx + col.offset());
col = col.child(lists_column_view::child_column_index);
auto list_col = cudf::detail::lists_column_device_view(col);
idx = list_col.offset_at(idx);
col = list_col.child();
}
}
return idx;
Expand Down
73 changes: 73 additions & 0 deletions cpp/tests/io/parquet_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3206,4 +3206,77 @@ TEST_F(ParquetWriterTest, RowGroupSizeInvalid)
cudf::logic_error);
}

TEST_F(ParquetWriterTest, EmptyList)
{
auto L1 = cudf::make_lists_column(0,
cudf::make_empty_column(cudf::data_type(cudf::type_id::INT32)),
cudf::make_empty_column(cudf::data_type{cudf::type_id::INT64}),
0,
{});
auto L0 = cudf::make_lists_column(
3, cudf::test::fixed_width_column_wrapper<int32_t>{0, 0, 0, 0}.release(), std::move(L1), 0, {});

auto filepath = temp_env->get_temp_filepath("EmptyList.parquet");
cudf::io::write_parquet(cudf::io::parquet_writer_options_builder(cudf::io::sink_info(filepath),
cudf::table_view({*L0})));

auto result = cudf_io::read_parquet(
cudf::io::parquet_reader_options_builder(cudf::io::source_info(filepath)));

using lcw = cudf::test::lists_column_wrapper<int64_t>;
auto expected = lcw{lcw{}, lcw{}, lcw{}};
cudf::test::print(expected);
cudf::test::expect_columns_equal(result.tbl->view().column(0), expected);
}

TEST_F(ParquetWriterTest, DeepEmptyList)
{
// Make a list column LLLi st only L is valid and LLi are all null. This tests whether we can
// handle multiple nullptr offsets

auto L2 = cudf::make_lists_column(0,
cudf::make_empty_column(cudf::data_type(cudf::type_id::INT32)),
cudf::make_empty_column(cudf::data_type{cudf::type_id::INT64}),
0,
{});
auto L1 = cudf::make_lists_column(
0, cudf::make_empty_column(cudf::data_type(cudf::type_id::INT32)), std::move(L2), 0, {});
auto L0 = cudf::make_lists_column(
3, cudf::test::fixed_width_column_wrapper<int32_t>{0, 0, 0, 0}.release(), std::move(L1), 0, {});

auto filepath = temp_env->get_temp_filepath("DeepEmptyList.parquet");
cudf::io::write_parquet(cudf::io::parquet_writer_options_builder(cudf::io::sink_info(filepath),
cudf::table_view({*L0})));

auto result = cudf_io::read_parquet(
cudf::io::parquet_reader_options_builder(cudf::io::source_info(filepath)));

cudf::test::expect_columns_equal(result.tbl->view().column(0), *L0);
}

TEST_F(ParquetWriterTest, EmptyListWithStruct)
{
auto L2 = cudf::make_lists_column(0,
cudf::make_empty_column(cudf::data_type(cudf::type_id::INT32)),
cudf::make_empty_column(cudf::data_type{cudf::type_id::INT64}),
0,
{});

auto children = std::vector<std::unique_ptr<cudf::column>>{};
children.push_back(std::move(L2));
auto S2 = cudf::make_structs_column(0, std::move(children), 0, {});
auto L1 = cudf::make_lists_column(
0, cudf::make_empty_column(cudf::data_type(cudf::type_id::INT32)), std::move(S2), 0, {});
auto L0 = cudf::make_lists_column(
3, cudf::test::fixed_width_column_wrapper<int32_t>{0, 0, 0, 0}.release(), std::move(L1), 0, {});

auto filepath = temp_env->get_temp_filepath("EmptyListWithStruct.parquet");
cudf::io::write_parquet(cudf::io::parquet_writer_options_builder(cudf::io::sink_info(filepath),
cudf::table_view({*L0})));
auto result = cudf_io::read_parquet(
cudf::io::parquet_reader_options_builder(cudf::io::source_info(filepath)));

cudf::test::expect_columns_equal(result.tbl->view().column(0), *L0);
}

CUDF_TEST_PROGRAM_MAIN()