Skip to content
Merged
Changes from all commits
Commits
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
30 changes: 11 additions & 19 deletions cpp/src/io/parquet/page_decode.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -25,16 +25,14 @@ struct page_state_s {
CUDF_HOST_DEVICE constexpr page_state_s() noexcept {}
uint8_t const* data_start{};
uint8_t const* data_end{};
uint8_t const* lvl_end{};
uint8_t const* dict_base{}; // ptr to dictionary page data
int32_t dict_size{}; // size of dictionary data
int32_t first_row{}; // First row in page to output
int32_t num_rows{}; // Rows in page to decode (including rows to be skipped)
int32_t first_output_value{}; // First value in page to output
int32_t num_input_values{}; // total # of input/level values in the page
int32_t dtype_len{}; // Output data type length
int32_t dtype_len_in{}; // Can be larger than dtype_len if truncating 32-bit into 8-bit
int32_t dict_bits{}; // # of bits to store dictionary indices
uint8_t const* dict_base{}; // ptr to dictionary page data
int32_t dict_size{}; // size of dictionary data
int32_t first_row{}; // First row in page to output
int32_t num_rows{}; // Rows in page to decode (including rows to be skipped)
int32_t num_input_values{}; // total # of input/level values in the page
int32_t dtype_len{}; // Output data type length
int32_t dtype_len_in{}; // Can be larger than dtype_len if truncating 32-bit into 8-bit
int32_t dict_bits{}; // # of bits to store dictionary indices
uint32_t dict_run{};
int32_t dict_val{};
uint32_t initial_rle_run[NUM_LEVEL_TYPES]{}; // [def,rep]
Expand All @@ -50,10 +48,9 @@ struct page_state_s {
int32_t ts_scale{}; // timestamp scale: <0: divide by -ts_scale, >0: multiply by ts_scale

// repetition/definition level decoding
int32_t input_value_count{}; // how many values of the input we've processed
int32_t input_row_count{}; // how many rows of the input we've processed
int32_t input_leaf_count{}; // how many leaf values of the input we've processed
uint8_t const* lvl_start[NUM_LEVEL_TYPES]{}; // [def,rep]
int32_t input_value_count{}; // how many values of the input we've processed
int32_t input_row_count{}; // how many rows of the input we've processed
int32_t input_leaf_count{}; // how many leaf values of the input we've processed
uint8_t const* abs_lvl_start[NUM_LEVEL_TYPES]{}; // [def,rep]
uint8_t const* abs_lvl_end[NUM_LEVEL_TYPES]{}; // [def,rep]
int32_t row_index_lower_bound{}; // lower bound of row indices we should process
Expand Down Expand Up @@ -983,7 +980,6 @@ inline __device__ uint32_t InitLevelSection(page_state_s* s,
s->initial_rle_value[lvl] = 0;
}
}
s->lvl_start[lvl] = cur;

if (cur > end) { s->set_error_code(decode_error::LEVEL_STREAM_OVERRUN); }
};
Expand All @@ -999,7 +995,6 @@ inline __device__ uint32_t InitLevelSection(page_state_s* s,
len = 0;
s->initial_rle_run[lvl] = s->page.num_input_values * 2; // repeated value
s->initial_rle_value[lvl] = 0;
s->lvl_start[lvl] = cur;
s->abs_lvl_start[lvl] = cur;
} else if (encoding == Encoding::RLE) { // V1 header with RLE encoding
if (cur + 4 < end) {
Expand All @@ -1017,7 +1012,6 @@ inline __device__ uint32_t InitLevelSection(page_state_s* s,
len = (s->page.num_input_values * level_bits + 7) >> 3;
s->initial_rle_run[lvl] = ((s->page.num_input_values + 7) >> 3) * 2 + 1; // literal run
s->initial_rle_value[lvl] = 0;
s->lvl_start[lvl] = cur;
s->abs_lvl_start[lvl] = cur;
} else {
len = 0;
Expand Down Expand Up @@ -1309,7 +1303,6 @@ inline __device__ bool setup_local_page_info(page_state_s* const s,
}
}
}
s->first_output_value = 0;

// Find the compressed size of repetition levels
cur += InitLevelSection(s, cur, end, level_type::REPETITION);
Expand Down Expand Up @@ -1369,7 +1362,6 @@ inline __device__ bool setup_local_page_info(page_state_s* const s,
}
}
if (cur > end) { s->set_error_code(decode_error::DATA_STREAM_OVERRUN); }
s->lvl_end = cur;
s->data_start = cur;
s->data_end = end;
} else {
Expand Down
Loading