From c03980503dc42f3c4fd6b95a9230e18d017d8737 Mon Sep 17 00:00:00 2001 From: Paul Mattione Date: Tue, 29 Oct 2024 18:17:31 -0400 Subject: [PATCH] fix merge issues --- cpp/src/io/parquet/decode_fixed.cu | 10 +++++----- cpp/src/io/parquet/parquet_gpu.hpp | 2 -- cpp/src/io/parquet/rle_stream.cuh | 4 ++-- 3 files changed, 7 insertions(+), 9 deletions(-) diff --git a/cpp/src/io/parquet/decode_fixed.cu b/cpp/src/io/parquet/decode_fixed.cu index 1ca2c404071..8d18c4c3e20 100644 --- a/cpp/src/io/parquet/decode_fixed.cu +++ b/cpp/src/io/parquet/decode_fixed.cu @@ -870,7 +870,7 @@ inline __device__ void bool_plain_decode(page_state_s* s, state_buf* sb, int t, if (t == 0) { s->dict_pos = pos; } } -template +template __device__ int skip_decode(stream_type& parquet_stream, int num_to_skip, int t) { // it could be that (e.g.) we skip 5000 but starting at row 4000 we have a run of length 2000: @@ -879,7 +879,7 @@ __device__ int skip_decode(stream_type& parquet_stream, int num_to_skip, int t) int num_skipped = parquet_stream.skip_decode(t, num_to_skip); while (num_skipped < num_to_skip) { // TODO: Instead of decoding, skip within the run to the appropriate location - auto const to_decode = min(2 * decode_block_size_t, num_to_skip - num_skipped); + auto const to_decode = min(rolling_buf_size, num_to_skip - num_skipped); num_skipped += parquet_stream.decode_next(t, to_decode); __syncthreads(); } @@ -1036,11 +1036,11 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size_t, 8) auto const skipped_leaf_values = s->page.skipped_leaf_values; if (skipped_leaf_values > 0) { if (should_process_nulls) { - skip_decode(def_decoder, skipped_leaf_values, t); + skip_decode(def_decoder, skipped_leaf_values, t); } - processed_count = skip_decode(rep_decoder, skipped_leaf_values, t); + processed_count = skip_decode(rep_decoder, skipped_leaf_values, t); if constexpr (has_dict_t) { - skip_decode(dict_stream, skipped_leaf_values, t); + skip_decode(dict_stream, skipped_leaf_values, t); } } } diff --git a/cpp/src/io/parquet/parquet_gpu.hpp b/cpp/src/io/parquet/parquet_gpu.hpp index 1c1185d8afd..8124be52552 100644 --- a/cpp/src/io/parquet/parquet_gpu.hpp +++ b/cpp/src/io/parquet/parquet_gpu.hpp @@ -994,7 +994,6 @@ void DecodePageDataFixedDict(cudf::detail::hostdevice_span pages, * @param[in] level_type_size Size in bytes of the type for level decoding * @param[in] has_nesting Whether or not the data contains nested (but not list) data. * @param[in] is_list Whether or not the data contains list data. - * @param[in] is_list Whether or not the data contains list data. * @param[out] error_code Error code for kernel failures * @param[in] stream CUDA stream to use */ @@ -1005,7 +1004,6 @@ void DecodeSplitPageFixedWidthData(cudf::detail::hostdevice_span pages int level_type_size, bool has_nesting, bool is_list, - bool is_list, kernel_error::pointer error_code, rmm::cuda_stream_view stream); diff --git a/cpp/src/io/parquet/rle_stream.cuh b/cpp/src/io/parquet/rle_stream.cuh index 76a2b050b9e..3c49de0c997 100644 --- a/cpp/src/io/parquet/rle_stream.cuh +++ b/cpp/src/io/parquet/rle_stream.cuh @@ -216,7 +216,7 @@ struct rle_stream { decode_index = -1; // signals the first iteration. Nothing to decode. } - __device__ inline int get_rle_run_info(rle_run& run) + __device__ inline int get_rle_run_info(rle_run& run) { run.start = cur; run.level_run = get_vlq32(run.start, end); @@ -383,7 +383,7 @@ struct rle_stream { // started basically we're setting up the rle_stream vars necessary to start fill_run_batch for // the first time while (cur < end) { - rle_run run; + rle_run run; int run_bytes = get_rle_run_info(run); if ((output_pos + run.size) > target_count) {