perf(parquet): decode def and rep levels concurrently via 2D grid - #7
Closed
vyasr wants to merge 10 commits into
Closed
perf(parquet): decode def and rep levels concurrently via 2D grid#7vyasr wants to merge 10 commits into
vyasr wants to merge 10 commits into
Conversation
…_page_decode_state replacing page_state_s (NVIDIA#23610) This PR completes the incremental narrowing pattern from NVIDIA#23471, NVIDIA#23479, NVIDIA#23495, and NVIDIA#23496 by retiring `page_state_s` entirely. It extracts `page_decode_nesting_state` (the nesting decode cache and pointer) and composes `full_page_decode_state` from `page_decode_setup_state`, `page_decode_stream_state`, `page_decode_nesting_state`, `page_decode_progress_state`, and `page_decode_output_state`. Every full-decode kernel (`decode_page_data`, `decode_split_page_data_kernel`, `decode_page_data_generic`, `decode_delta_binary`, `decode_delta_byte_array`, `decode_delta_length_byte_array`, `compute_string_page_bounds`, `compute_page_sizes`) now takes this composed state instead of the ad-hoc `page_state_s`, and the old struct is deleted. The composition also folds `page_state_s`'s stand-alone level-decoding fields (`lvl_start[2]`, `lvl_end`) into the already-existing `page_decode_stream_state::abs_lvl_start` / `abs_lvl_end` arrays, and drops the unused `first_output_value` field. Together with the composition's tighter layout this removes 32 bytes of shared memory from every full-decode kernel on every generated architecture, with no register-count regressions. A fresh `cuobjdump -res-usage` comparison against `upstream/main` shows the same 32-byte shmem reduction on all eight migrated kernels on every generated architecture. Register usage does not increase anywhere; one variant of `decode_split_page_data_kernel` on `sm_86` drops by 8 registers: | Kernel | Arch summary | shmem delta | register changes | |---|---|---:|---| | `decode_page_data` | all generated arch variants | -32 bytes | unchanged | | `decode_split_page_data_kernel` | all generated arch variants | -32 bytes | unchanged except `sm_86`: one template variant 48 -> 40 | | `decode_page_data_generic` | all generated arch variants (44 template variants) | -32 bytes | unchanged | | `compute_page_sizes` | all generated arch variants | -32 bytes | unchanged | | `compute_string_page_bounds` | all generated arch variants | -32 bytes | unchanged | | `decode_delta_binary` | all generated arch variants | -32 bytes | unchanged | | `decode_delta_byte_array` | all generated arch variants | -32 bytes | unchanged | | `decode_delta_length_byte_array` | all generated arch variants | -32 bytes | unchanged | Representative sm_80 measurements: | Kernel | shmem before | shmem after | delta | regs before | regs after | |---|---:|---:|---:|---:|---:| | `decode_page_data` | 4112 | 4080 | -32 | 56 | 56 | | `decode_split_page_data_kernel` | 4112 | 4080 | -32 | 56 | 56 | | `decode_page_data_generic` (min variant) | 2104 | 2072 | -32 | 62 | 62 | | `decode_page_data_generic` (max variant) | 13024 | 12992 | -32 | 64 | 64 | | `compute_page_sizes` | 3240 | 3208 | -32 | 32 | 32 | | `compute_string_page_bounds` | 3256 | 3224 | -32 | 32 | 32 | | `decode_delta_binary` | 2700 | 2668 | -32 | 56 | 56 | | `decode_delta_byte_array` | 5068 | 5036 | -32 | 72 | 72 | | `decode_delta_length_byte_array` | 3392 | 3360 | -32 | 64 | 64 | With this PR the `page_state_s` type is fully removed. All Parquet decode and preprocess kernels now use one of four purpose-built shared-memory states (`level_scan_state`, `string_size_scan_state`, `string_offset_scan_state`, `full_page_decode_state`), each holding only the substructs it actually needs. Authors: - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - Bradley Dice (https://github.com/bdice) URL: NVIDIA#23610
…ck (NVIDIA#23314) The GPU Parquet reader rejects DELTA-encoded pages whose mini-blocks hold more than 64 values with `DELTA_PARAMS_UNSUPPORTED`, although the format allows any multiple of 32 and other readers accept such files. This PR removes the mini-block size limit from the DELTA_BINARY_PACKED, DELTA_BYTE_ARRAY and DELTA_LENGTH_BYTE_ARRAY decoders. The limit came from the decoder's rolling value buffer, which had to hold two whole mini-blocks. Instead of growing it, decode and consume mini-blocks one warp-size pass at a time, so every buffer is sized by the decode pipeline (values in flight per iteration) rather than by any mini-block geometry: - The decode kernels' producer warps decode one 32-value pass per call (`delta_binary_decoder::decode_next_pass()`); each page still produces `min(values_per_mb, 64)` values per main-loop iteration, so previously-readable pages keep their exact iteration schedule. Shared memory for the value buffers is unchanged. - The skip paths (`skip_values`, `skip_values_and_sum`, `delta_byte_array_decoder::skip`) also advance pass by pass and resume at a pass boundary, instead of requiring a whole mini-block to stay resident; pages resuming after a skip produce one pass per iteration since larger batches could overwrite the up-to-31 not-yet-consumed values the skip leaves behind. - The string-size prepasses read decoded lengths back per pass; whole-mini-block read-back silently mis-computes `str_bytes` once a mini-block no longer fits the buffer. - `init_binary_block` now validates `values_per_mb % 32 == 0`; malformed headers previously decoded garbage instead of erroring. This also fixes two latent bugs reachable on nested pages with the old code: - The `nz_idx` ring (leaf-ordinal to output-row map) shares its size with the value buffers, but on nested pages the level decoder overshoots its target by up to a warp of values and wraps onto entries the value consumer is reading when a page decodes 64 values per iteration (`values_per_mb == 64`, e.g. arrow-rs INT64 lists; confirmed with compute-sanitizer racecheck). The ring now has its own, larger size. - `delta_byte_array_decoder::skip` saved the string needed for the next batch's front compression at a stale offset inside the scratch it was about to overwrite; the last decoded string now lives in a reserved slot past the scratch area. Since no stock writer emits more than 64 values per mini-block (cudf and parquet-mr write 32, pyarrow and arrow-rs write 64 for INT64), tests build single-page files in memory from plain value vectors with new test utilities (`parquet_delta_test_utils.hpp`; the builders' output was cross-checked against pyarrow for every geometry the tests use): flat INT64, LIST<INT64>, and both string encodings flat and as LIST<STRING>, at 64/96/128/256 values per mini-block, with full, `num_rows`-trimmed and `skip_rows` reads. Also adds `parquet_read_delta_binary`/`parquet_read_delta_string` benchmarks (the DELTA decode kernels had no reader benchmark; results are within ~1% of the previous code for DELTA_BINARY_PACKED and DELTA_BYTE_ARRAY, ~2% for DELTA_LENGTH_BYTE_ARRAY). Authors: - Pramod Satya (https://github.com/pramodsatya) - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - Paul Mattione (https://github.com/pmattione-nvidia) - Muhammad Haseeb (https://github.com/mhaseeb123) URL: NVIDIA#23314
Contributes to NVIDIA#23464 `encode_columns` allocates one `device_uvector` per (stripe, stream) pair, so a wide table with many stripes produces thousands of device allocations. With this PR, we compute the size of every stream up front and suballocate an extent (an aligned byte range within an arena) for each. Each extent is aligned exactly as a fresh RMM allocation would be. Extent sizes are upper bounds, because the encoded length of an RLE stream is not known until it has been encoded. `gather_stripes` compacts a stripe's per-rowgroup chunks into a contiguous, tightly sized stream. The encoder output is therefore split across two arenas, by whether an extent might be compacted: - **`transient_buffer`** - extents whose stripe spans several rowgroups *and* whose size is a strict upper bound rather than exact. `gather_stripes` compacts every extent in this arena, which lets the whole arena be released as soon as gathering completes, before compression allocates. - **`persistent_buffer`** - everything else: single-rowgroup stripes, and multi-rowgroup extents whose size estimate is exact. This arena lives until the encoded data is no longer needed. - **`gathered_buffer`** - the compaction destination. Authors: - Vukasin Milovanovic (https://github.com/vuule) Approvers: - Muhammad Haseeb (https://github.com/mhaseeb123) - Tianyu Liu (https://github.com/kingcrimsontianyu) URL: NVIDIA#23433
…ns (NVIDIA#23564) Doxygen tags in `cpp/include` and `cpp/src` that name something the declaration below them does not have. Comments only — no code, no behaviour, no tests. **`@tparam` with the name missing** — `column_view_base::head` and `mutable_column_view::head` both carry `@tparam The type to cast to`, so Doxygen takes `The` as the template parameter name and `T` goes undocumented. **`@tparam` left over from a refactor** — eight of the `make_device_uvector*` overloads in `detail/utilities/vector_factories.hpp` document both `Container` and `T`, but those overloads are `template <typename Container>` only. The `T` line is a leftover from the `host_span<T>` overloads above them. **`@tparam` renamed** — `id_to_type` documents `t` while the parameter is `Id`; `dictionary_column_wrapper` documents `SourceElementTo` while the parameter is `SourceElementT`. **`@param` on a function that takes nothing** — `arrow_column::view()` and `arrow_table::view()` document `stream` and `mr`, copied from the `to_arrow` overload directly above each. Both take no arguments. **`@param` where `@tparam` was meant** — `is_boolean()` documents `@param type The data_type to verify` and `is_nested()` documents `@param T`; both are `template <typename T>` predicates with no runtime arguments. A second commit covers the further files pointed out in review: `column_device_view_base.cuh`, `type_dispatcher.hpp`'s `IdTypeMap`, `iterator.cuh`, `timestamp_utilities.cuh`, `rolling/detail/range_utils.cuh` and `strings/count_matches.hpp`. Follows NVIDIA#21762 and NVIDIA#21764, which did the same for other entries in `/include` and `/src`. Every tag here was opened and read against the declaration underneath it. Authors: - Dmitry (https://github.com/darkdi) - Igor Peshansky (https://github.com/igorpeshansky) - Nghia Truong (https://github.com/ttnghia) Approvers: - Igor Peshansky (https://github.com/igorpeshansky) - Nghia Truong (https://github.com/ttnghia) - Yunsong Wang (https://github.com/PointKernel) URL: NVIDIA#23564
…3625) XREF: rapidsai/build-infra#374 CODEOWNER teams must be in the same org as the repo. Let's update them now that we've migrated to the NVIDIA org.
…#23271) This PR implements an alternative approach for the RLE decoding. The old approach used a producer-consumer model where warp 0 populates a ring buffer of runs for other warps to pick off. That leads to two sources of imbalance: 1. Warp 0 becomes a bottleneck for the other warps because production can't keep up with decode. 2. Different warps operate on runs of different lengths, leading to interwarp imbalances even among consumer warps. With the new approach, the full stream is split into chunks of a fixed size (determined at compile time). Within each chunk, thread 0 does a serial pass through the data to find all of the RLE headers and populates the associated splits and metadata (RLE vs bit-packed) into a shared memory array. Then, all warps can cooperatively read through all of that data. Since data is parsed by chunk rather than by run, there is no longer any imbalance between warps. Warps keep track of boundaries via the same shared memory arrays, and therefore warps can start in the middle of any run and completely traverse runs short enough to fit within their chunks. Authors: - Vyas Ramasubramani (https://github.com/vyasr) - Muhammad Haseeb (https://github.com/mhaseeb123) Approvers: - Muhammad Haseeb (https://github.com/mhaseeb123) - Paul Mattione (https://github.com/pmattione-nvidia) URL: NVIDIA#23271
Fixes some code logic where strings offsets were hardcoded to expect int32 (or cudf::size_type). These were found while investigating size_type/offset-type usage in libcudf. Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Bradley Dice (https://github.com/bdice) - Lawrence Mitchell (https://github.com/wence-) URL: NVIDIA#23612
Each page's level-decode kernel is now dispatched with a 2D grid: dim_grid = (pages.size(), 2) blockIdx.y selects the level stream (DEFINITION=0, REPETITION=1). Blocks for absent streams (non-null pages, non-list pages) return immediately, so occupancy is preserved. The two streams decode concurrently on separate SMs instead of sequentially within one block.
Assert that the level-stream range is non-negative on entry to rle_stream::init(). Catches callers that pass an inverted range early rather than silently producing garbage output. Suggested by reviewer r3760481306.
vyasr
force-pushed
the
opt/rle-def-rep-split
branch
from
August 12, 2026 19:15
9e67823 to
033720a
Compare
Owner
Author
|
Upstreamed in NVIDIA#23637. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
6 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Switch the kernel from a 1D to a 2D grid where the y axis determines whether to decode definition or repetition levels. That reduces some synchronization and local state.