Optimize RLE decoding using a warp-balanced chunking approach - #23271
Optimize RLE decoding using a warp-balanced chunking approach#23271vyasr wants to merge 37 commits into
Conversation
vyasr
added
Performance
…chunked expand Readability-only. Ptxas SMEM unchanged at 5168 B for preprocess_levels_kernel on SM 8.0. All ctest -R PARQUET pass including PARQUET_RLE_CHUNKED_EQUIVALENCE_TEST. No performance change expected.
Dead scaffolding from earlier iterations of the chunked-expand path. The actual decode_next_chunked function allocates gen_out_off and gen_meta directly as __shared__ arrays; this struct was never instantiated anywhere in the tree. No functional change. SMEM/register footprint of preprocess_levels_kernel unchanged (5168 B / 48 reg on SM 8.0). All PARQUET tests pass.
Both decode_next_ring and decode_next_chunked had the same all-zeros short-circuit for level_bits == 0. Pull it up into decode_next so the helpers stay focused on the general RLE path. Uses cur_values-relative addressing (the ring version's form was only correct because cur_values was always 0 on entry in practice).
…nce test The rle_stream::init signature added a Group parameter as part of the SMEM-staging refactor; update the chunked-equivalence test to pass cg::this_thread_block() to match.
A100 nvbench parquet_read_decode sweep on LIST/STRUCT/STRING shows k=1024 is 80-95% faster than k=512 on the chunked-expand path in preprocess_levels_kernel. k=2048 saturates or regresses ~10% on LIST due to occupancy pressure, so 1024 is the sweet spot. Keep k=512 on sm_70 (V100) where the larger SMEM footprint would exceed the preprocess_levels_kernel budget. H100/Blackwell tiers are left as TODO in the comment; sm_80+ tier applies until they are measured. SMEM cost: (2 * kGenRuns + 1) * 4 bytes = 8196 B at k=1024 (vs 4100 B at k=512). Well within A100/H100 budgets.
level_mask is invariant across the whole decode_next_chunked call (depends only on level_bits, which is class state). Move it out to the top of the function. Verified with cuobjdump --dump-resource-usage that this is a no-op for register counts on sm_80/86/90 - the compiler was already hoisting it - so this is a readability/intent-signalling change only.
The earlier sm_80+ bump to 1024 was justified by an isolated kGenRuns sweep whose k=512 STRUCT baseline was in a pathological regime (245-415 ms with 46-75% noise, hitting nvbench timeouts). A full 36-config parquet_read_decode A/B against upstream/main on A100 shows that k=1024 vs k=512 differences are within noise, and a matching sweep on H100 found the same pattern - k=2048 was numerically best but within noise of both k=512 and k=1024. Since no architecture shows a meaningful preference in real workloads, drop the arch-adaptive split and use a single kGenRuns=512 everywhere. This also reclaims 4 KiB of SMEM per block on sm_80+, easing occupancy pressure on register-constrained architectures like sm_86.
Rename the identifiers we introduced with the chunked-expand path to
match the snake_case convention used throughout libcudf:
kGenRuns -> max_runs_per_chunk (also better describes what the
constant controls)
kWarps -> num_warps (matches existing num_rle_stream_*
naming)
gen_out_off, gen_meta (and their _v span views) -> chunk_out_off,
chunk_meta (the opaque "gen_" prefix came from
"generated in phase 1"; "chunk_"
parallels the existing s_chunk_*
shared vars produced alongside them)
Verified with build-cudf-cpp and PARQUET_RLE_CHUNKED_EQUIVALENCE_TEST (8/8).
warp_fill was only used from the RLE-run arm of the chunked-expand loop. Inline it at the call site and drop the helper. This removes the now-unnecessary __forceinline__ / __restrict__ annotations and the comments that referenced warp_fill by name; the loop is short enough to be self-explanatory at the call site.
Sweeps of {256, 512, 1024, 2048, 4096} on A100, H100, and B200 all show
1024 either as the numerical optimum or within noise of it (H100 leaned
toward 2048 by a small margin, still within noise). The delta over 512
is small - typically a few percent - but 1024 is the consistent winner
across the modern arches, so specialize by __CUDA_ARCH__.
sm_70/sm_75 stay at 512 because 1024 does not fit the
preprocess_levels_kernel SMEM budget on those older architectures.
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
📝 WalkthroughWalkthroughThe Parquet RLE stream gains an optional chunked-expand decoder with partial-run resumption and configurable staging capacity. Parquet level preprocessing selects this mode and supplies the staging size during repetition- and definition-level decoder initialization. ChangesParquet RLE chunked decoding
Estimated code review effort: 4 (Complex) | ~45 minutes Possibly related PRs
Suggested labels: Suggested reviewers: 🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
🧹 Nitpick comments (2)
cpp/tests/io/parquet_rle_chunked_equivalence_test.cu (2)
6-20: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winMissing explicit gtest include.
TEST_F/EXPECT_EQare used directly but the file doesn't include<cudf_test/cudf_gtest.hpp>(or rawgtest/gtest.h); it relies on a transitive include viacudf_test/base_fixture.hpp.As per coding guidelines, "Test files must include
#include <cudf_test/cudf_gtest.hpp>instead of rawgtest/gtest.h."🧪 Proposed fix
`#include` "../../src/io/parquet/rle_stream.cuh" +#include <cudf_test/cudf_gtest.hpp> `#include` <cudf_test/base_fixture.hpp> `#include` <cudf_test/testing_main.hpp>🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/tests/io/parquet_rle_chunked_equivalence_test.cu` around lines 6 - 20, Add the explicit <cudf_test/cudf_gtest.hpp> include to parquet_rle_chunked_equivalence_test.cu alongside the other cudf_test headers, so TEST_F and EXPECT_EQ do not rely on the transitive inclusion from base_fixture.hpp.Source: Coding guidelines
179-186: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win
ManyShortRepeatedRunsdoesn't reliably cross the chunk boundary on sm_80+.600 short runs is below
max_runs_per_chunk(1024) on sm_80+, so the outer multi-chunk loop indecode_next_chunked(which is the core new logic this PR adds) may not actually be exercised on the modern architectures this feature targets, only on sm_70/75 (max 512). Consider sizing the run count offcudf::io::parquet::detail::max_runs_per_chunkdirectly (e.g.max_runs_per_chunk * 2 + 1) so the test deterministically crosses the boundary regardless of the GPU running CI.As per coding guidelines, "Test suites should cover edge cases such as empty input, null values, sliced columns, boundary sizes, and multi-block sizes."
TEST_F(ParquetRleChunkedEquivalenceTest, ManyShortRepeatedRuns) { using cudf::io::parquet::detail::max_runs_per_chunk; int const num_runs = max_runs_per_chunk * 2 + 1; // guarantees >1 chunk on any arch std::vector<uint8_t> encoded; for (int i = 0; i < num_runs; ++i) { append_repeated(encoded, 1, i & 15, 4); } run_case<uint8_t>(encoded, 4, num_runs, num_runs); }🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/tests/io/parquet_rle_chunked_equivalence_test.cu` around lines 179 - 186, Update the ManyShortRepeatedRuns test to derive its run count from cudf::io::parquet::detail::max_runs_per_chunk, using a value greater than two chunk capacities (such as 2 * max_runs_per_chunk + 1). Use that count in the encoding loop and run_case expectations so the test deterministically exercises decode_next_chunked’s multi-chunk path on every supported architecture.Source: Coding guidelines
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Nitpick comments:
In `@cpp/tests/io/parquet_rle_chunked_equivalence_test.cu`:
- Around line 6-20: Add the explicit <cudf_test/cudf_gtest.hpp> include to
parquet_rle_chunked_equivalence_test.cu alongside the other cudf_test headers,
so TEST_F and EXPECT_EQ do not rely on the transitive inclusion from
base_fixture.hpp.
- Around line 179-186: Update the ManyShortRepeatedRuns test to derive its run
count from cudf::io::parquet::detail::max_runs_per_chunk, using a value greater
than two chunk capacities (such as 2 * max_runs_per_chunk + 1). Use that count
in the encoding loop and run_case expectations so the test deterministically
exercises decode_next_chunked’s multi-chunk path on every supported
architecture.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 0c826ec9-5852-4d01-9114-574ad1224e6b
📒 Files selected for processing (4)
cpp/src/io/parquet/decode_preprocess.cucpp/src/io/parquet/rle_stream.cuhcpp/tests/CMakeLists.txtcpp/tests/io/parquet_rle_chunked_equivalence_test.cu
Applies the mechanical batch of review comments from PR rapidsai#23271: - Remove stale `// otherwise, full decode.` comment in decode_next_ring (leftover from a removed branch). - Replace hand-written `(x + N - 1) / N` with cudf::util::div_rounding_up_{safe,unsafe} in decode_next_chunked (level_bits->byte width, per-warp chunk slice size). - Rename single-letter locals in decode_next_chunked for readability: Phase 1: n -> num_runs, co -> run_prefix_end, cnt -> run_len, meta -> run_desc Phase 2: a -> first_run_idx, r -> run_idx, r_lo/r_hi -> run_start_out/run_end_out, seg_lo/seg_hi -> out_lo/out_hi, meta -> run_desc Comments referring to those names updated accordingly. - Add missing `#include <cudf_test/cudf_gtest.hpp>` in the RLE chunked equivalence test. No functional changes. PARQUET_RLE_CHUNKED_EQUIVALENCE_TEST passes (8/8), and TPC-H SF1000 sum-of-min iteration times are within noise of the pre-change baseline.
Per PR rapidsai#23271 review (mhaseeb123): the existing parquet reader test suite (PARQUET_TEST, PARQUET_DELETION_VECTORS_TEST, STREAM_IO_PARQUET_TEST) already exercises both the ring-buffer and chunked-expand rle_stream paths through real parquet page data, so the standalone equivalence harness is redundant. Removing it eliminates 213 lines of unit-test scaffolding that duplicates existing coverage. The remaining parquet tests continue to exercise the chunked-expand path via the use_chunked_expand=true rle_stream instantiation used during page decode.
Addresses PR rapidsai#23271 review comments #3, #4, rapidsai#11 from mhaseeb123: - Template `decode_next_chunked` on a cooperative_groups `Group` type, matching the pattern already used by `rle_stream::init`. The public entrypoint `decode_next(int t, int count)` forwards to `decode_next_chunked(cg::this_thread_block(), count)` so no call sites in decode_fixed.cu / decode_preprocess.cu change. - Replace hand-computed lane and warp indexing with cg primitives: t & 31 -> warp.thread_rank() t >> 5 -> warp.meta_group_rank() (renamed local: warp_id) hardcoded 32 -> warp.size() num_rle_stream_decode_threads / warp_size -> warp.meta_group_size() where `warp = cg::tiled_partition<cudf::detail::warp_size>(group)`. - Replace the two `__syncthreads()` calls with `group.sync()`. - Replace `if (t == 0) { ... }` with `cg::invoke_one(group, [&](){ ... });` and refresh the surrounding comment to match. - Add a `namespace cg = cooperative_groups;` alias inside cudf::io::parquet::detail, following the convention already used in decode_preprocess.cu. Codegen is intended to be a no-op: `cg::tiled_partition<32>(block)` and its accessors are all `_CG_STATIC_QUALIFIER` inline, `group.sync()` on a `thread_block` lowers to `__syncthreads()`, and `cg::invoke_one` on a block lowers to `if (group.thread_rank() == 0) fn()` (no implicit sync) per <cooperative_groups/details/invoke.h>. Comparing object files built pre- and post-refactor confirms this: decode_fixed.cu.o 9,715,392 -> 9,715,392 (identical) page_string_decode.cu.o 1,768,912 -> 1,768,912 (identical) decode_preprocess.cu.o 807,048 -> 812,848 (+0.7% debug info) Verification: - PARQUET_TEST 481/481 pass (1 pre-existing skip) - PARQUET_DELETION_VECTORS_TEST 6/6 pass - STREAM_IO_PARQUET_TEST 4/4 pass - TPC-H SF1000 sum-of-min: 221.09s vs 222.46s pre-refactor baseline (-0.6%, within run-to-run noise).
Addresses PR rapidsai#23271 review comment #1 from mhaseeb123: mirror the cooperative-groups treatment that commit 041c4ab applied to decode_next_chunked in the pre-existing ring path. - Template `decode_next_ring` on a cooperative_groups `Group` type and take `Group const& group` instead of `int t`. The dispatcher `decode_next(int t, int count)` still forwards from a raw thread id by calling `decode_next_ring(cg::this_thread_block(), count)`, so call sites in decode_fixed.cu / decode_preprocess.cu are untouched. - Replace hand-computed indexing with cg primitives: t / warp_size -> warp.meta_group_rank() (warp_id) t % warp_size -> warp.thread_rank() (warp_lane) where `warp = cg::tiled_partition<cudf::detail::warp_size>(group)`. - Replace the three `__syncthreads()` calls with `group.sync()` and the single `__syncwarp()` with `warp.sync()`. - Replace the outer `if (t == 0) { ... }` block-level shared-var initialisation with `cg::invoke_one(group, [&](){ ... })`. - Replace the two `if (warp_lane == 0) { ... }` warp-level leader blocks with `cg::invoke_one(warp, [&](){ ... })`. As with the chunked variant, cg primitives on a `thread_block` lower to the same instructions as the original hand-coded versions: `warp.thread_rank()` -> `threadIdx.x & 31`, `group.sync()` -> `__syncthreads()`, `cg::invoke_one` on a block/tile -> `if (group.thread_rank() == 0) fn()` (no implicit sync, verified in <cooperative_groups/details/invoke.h>). libcudf.so grows by 56 KiB (0.005%) due to the additional lambda debug info; the code section is functionally unchanged. Verification: - PARQUET_TEST 481/481 pass (1 pre-existing skip) - PARQUET_DELETION_VECTORS_TEST 6/6 pass - STREAM_IO_PARQUET_TEST 4/4 pass - TPC-H SF1000 sum-of-min: 213.44s vs 222.46s pre-refactor baseline (-4.1%, i.e. no regression; the ring path is exercised by decode_fixed.cu dict/bool streams and by decode_preprocess.cu repetition-level decoding, so every query in the benchmark touches this code).
Rewrite the phase-2 expand loop so each lane iterates over its own output positions and walks run_idx forward, instead of the warp looping over runs. Per pmattione review comment on PR rapidsai#23271. Design: - Each lane owns positions p = lo + lane + k * warp.size(). - run_idx is set once via binary search on the lane's initial p, then advanced by linear walk (usually 0 steps when still in the same run). - The old per-run inner loop left most of the warp idle on short runs; this version keeps all 32 lanes writing every iteration. Amortization: - Total run-advances by all 32 lanes over the slice equals num_runs_in_slice (each boundary crossed once). Divided by 32 lanes in parallel, this is (num_runs_in_slice / 32) warp cycles of linear-walk cost across the whole slice - a rounding error against payload work. - Long-run pages: zero-step inner walks, same access pattern as before. No regression. TPC-H SF1000 A/B against the prior HEAD (post-cg baseline 213.440s sum-of-min): - New: 208.841s sum-of-min (-2.15%). - Big wins on dictionary-heavy queries: Q14 -11.5%, Q16 -12.6%, Q19 -12.5%, Q13 -11.2%. - Q18 +5.06% and Q21 +4.39% are the expected small tradeoff on long-run-heavy lineitem aggregations, where per-position setup is paid on every iteration instead of once per run. Net effect is still positive because the short-run wins dominate. Parquet gtests: PARQUET_TEST 481/481 (+1 pre-existing skip), PARQUET_DELETION_VECTORS_TEST 6/6, STREAM_IO_PARQUET_TEST 4/4.
Ultraworked with [Sisyphus](https://github.com/code-yeongyu/oh-my-openagent) Co-authored-by: Sisyphus <clio-agent@sisyphuslabs.ai>
|
On B200 the chunked-expand path is faster on all the microbenchmarks I tested with a geomean speedup of -2.4% across the full suite. The largest wins are, as expected, on the Full
|
| T | io_type | cardinality | run_length | chunk_read_limit | pass_read_limit | data_size | row_group_size_bytes | row_group_size_rows | Ref Time | Ref Noise | Cmp Time | Cmp Noise | Diff | %Diff | Status |
|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
| LIST | DEVICE_BUFFER | 0 | 1 | 0 | 0 | 536870912 | 0 | 0 | 12.935 ms | 0.43% | 11.769 ms | 0.23% | -1165.956 us | -9.01% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 1 | 0 | 0 | 536870912 | 0 | 0 | 14.310 ms | 0.29% | 14.188 ms | 0.28% | -122.651 us | -0.86% | FAST |
| LIST | DEVICE_BUFFER | 0 | 32 | 0 | 0 | 536870912 | 0 | 0 | 12.732 ms | 0.29% | 12.549 ms | 0.21% | -182.733 us | -1.44% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 32 | 0 | 0 | 536870912 | 0 | 0 | 12.459 ms | 0.25% | 12.327 ms | 0.21% | -131.886 us | -1.06% | FAST |
| LIST | DEVICE_BUFFER | 0 | 1 | 500000 | 0 | 536870912 | 0 | 0 | 177.615 ms | 0.46% | 174.849 ms | 0.02% | -2766.019 us | -1.56% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 1 | 500000 | 0 | 536870912 | 0 | 0 | 159.227 ms | 0.37% | 157.525 ms | 0.16% | -1702.182 us | -1.07% | FAST |
| LIST | DEVICE_BUFFER | 0 | 32 | 500000 | 0 | 536870912 | 0 | 0 | 168.299 ms | 0.19% | 167.429 ms | 0.18% | -869.760 us | -0.52% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 32 | 500000 | 0 | 536870912 | 0 | 0 | 166.106 ms | 1.59% | 165.119 ms | 0.13% | -986.718 us | -0.59% | FAST |
| LIST | DEVICE_BUFFER | 0 | 1 | 0 | 500000 | 536870912 | 0 | 0 | 410.169 ms | 0.04% | 405.354 ms | 0.16% | -4814.972 us | -1.17% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 1 | 0 | 500000 | 536870912 | 0 | 0 | 354.901 ms | 0.04% | 353.751 ms | 0.09% | -1150.616 us | -0.32% | FAST |
| LIST | DEVICE_BUFFER | 0 | 32 | 0 | 500000 | 536870912 | 0 | 0 | 229.962 ms | 0.01% | 224.759 ms | 0.03% | -5202.841 us | -2.26% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 32 | 0 | 500000 | 536870912 | 0 | 0 | 252.426 ms | 0.04% | 246.873 ms | 0.19% | -5552.954 us | -2.20% | FAST |
| LIST | DEVICE_BUFFER | 0 | 1 | 500000 | 500000 | 536870912 | 0 | 0 | 533.773 ms | 0.06% | 531.347 ms | 0.03% | -2425.964 us | -0.45% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 1 | 500000 | 500000 | 536870912 | 0 | 0 | 443.766 ms | 0.07% | 444.835 ms | 0.07% | 1.069 ms | 0.24% | SLOW |
| LIST | DEVICE_BUFFER | 0 | 32 | 500000 | 500000 | 536870912 | 0 | 0 | 343.355 ms | 0.03% | 340.103 ms | 0.12% | -3251.587 us | -0.95% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 32 | 500000 | 500000 | 536870912 | 0 | 0 | 362.383 ms | 0.05% | 358.248 ms | 0.08% | -4134.705 us | -1.14% | FAST |
parquet_read_fixed_width_struct
[0] NVIDIA B200
| data_type | io_type | compression_type | cardinality | run_length | data_size | row_group_size_bytes | row_group_size_rows | Ref Time | Ref Noise | Cmp Time | Cmp Noise | Diff | %Diff | Status |
|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
| STRUCT | DEVICE_BUFFER | NONE | 0 | 1 | 536870912 | 0 | 0 | 6.019 ms | 0.40% | 5.715 ms | 0.50% | -304.312 us | -5.06% | FAST |
| STRUCT | DEVICE_BUFFER | NONE | 1000 | 1 | 536870912 | 0 | 0 | 7.321 ms | 0.50% | 6.907 ms | 0.30% | -413.992 us | -5.66% | FAST |
| STRUCT | DEVICE_BUFFER | NONE | 0 | 32 | 536870912 | 0 | 0 | 6.564 ms | 0.57% | 6.270 ms | 0.37% | -293.521 us | -4.47% | FAST |
| STRUCT | DEVICE_BUFFER | NONE | 1000 | 32 | 536870912 | 0 | 0 | 6.592 ms | 0.66% | 6.286 ms | 0.79% | -305.872 us | -4.64% | FAST |
parquet_read_decode
[0] NVIDIA B200
| data_type | io_type | compression_type | cardinality | run_length | data_size | row_group_size_bytes | row_group_size_rows | Ref Time | Ref Noise | Cmp Time | Cmp Noise | Diff | %Diff | Status |
|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
| LIST | DEVICE_BUFFER | NONE | 0 | 1 | 536870912 | 0 | 0 | 10.176 ms | 0.27% | 9.053 ms | 0.17% | -1123.714 us | -11.04% | FAST |
| LIST | DEVICE_BUFFER | NONE | 1000 | 1 | 536870912 | 0 | 0 | 12.403 ms | 0.21% | 12.399 ms | 0.46% | -4.368 us | -0.04% | SAME |
| LIST | DEVICE_BUFFER | NONE | 0 | 32 | 536870912 | 0 | 0 | 11.302 ms | 0.19% | 11.231 ms | 0.48% | -70.714 us | -0.63% | FAST |
| LIST | DEVICE_BUFFER | NONE | 1000 | 32 | 536870912 | 0 | 0 | 11.183 ms | 0.21% | 11.072 ms | 0.23% | -111.383 us | -1.00% | FAST |
parquet_read_chunks
[0] NVIDIA B200
| T | io_type | cardinality | run_length | chunk_read_limit | data_size | row_group_size_bytes | row_group_size_rows | Ref Time | Ref Noise | Cmp Time | Cmp Noise | Diff | %Diff | Status |
|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
| LIST | DEVICE_BUFFER | 0 | 1 | 0 | 536870912 | 0 | 0 | 12.931 ms | 0.39% | 11.764 ms | 0.22% | -1166.752 us | -9.02% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 1 | 0 | 536870912 | 0 | 0 | 14.312 ms | 0.98% | 14.199 ms | 0.31% | -112.379 us | -0.79% | FAST |
| LIST | DEVICE_BUFFER | 0 | 32 | 0 | 536870912 | 0 | 0 | 12.692 ms | 0.19% | 12.555 ms | 0.24% | -136.163 us | -1.07% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 32 | 0 | 536870912 | 0 | 0 | 12.491 ms | 0.19% | 12.299 ms | 0.16% | -191.536 us | -1.53% | FAST |
| LIST | DEVICE_BUFFER | 0 | 1 | 500000 | 536870912 | 0 | 0 | 536.629 ms | 0.09% | 526.314 ms | 0.14% | -10315.759 us | -1.92% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 1 | 500000 | 536870912 | 0 | 0 | 422.651 ms | 0.04% | 419.647 ms | 0.08% | -3003.784 us | -0.71% | FAST |
| LIST | DEVICE_BUFFER | 0 | 32 | 500000 | 536870912 | 0 | 0 | 345.866 ms | 0.03% | 337.628 ms | 0.07% | -8238.361 us | -2.38% | FAST |
| LIST | DEVICE_BUFFER | 1000 | 32 | 500000 | 536870912 | 0 | 0 | 365.392 ms | 0.04% | 357.786 ms | 0.13% | -7606.744 us | -2.08% | FAST |
| STRUCT | DEVICE_BUFFER | 0 | 1 | 0 | 536870912 | 0 | 0 | 14.705 ms | 0.36% | 14.381 ms | 0.50% | -324.074 us | -2.20% | FAST |
| STRUCT | DEVICE_BUFFER | 1000 | 1 | 0 | 536870912 | 0 | 0 | 10.956 ms | 0.75% | 10.598 ms | 0.78% | -358.173 us | -3.27% | FAST |
| STRUCT | DEVICE_BUFFER | 0 | 32 | 0 | 536870912 | 0 | 0 | 17.369 ms | 0.40% | 17.117 ms | 0.50% | -251.305 us | -1.45% | FAST |
| STRUCT | DEVICE_BUFFER | 1000 | 32 | 0 | 536870912 | 0 | 0 | 10.936 ms | 0.74% | 10.609 ms | 0.50% | -327.587 us | -3.00% | FAST |
| STRUCT | DEVICE_BUFFER | 0 | 1 | 500000 | 536870912 | 0 | 0 | 158.617 ms | 0.14% | 158.957 ms | 0.36% | 339.655 us | 0.21% | SLOW |
| STRUCT | DEVICE_BUFFER | 1000 | 1 | 500000 | 536870912 | 0 | 0 | 115.347 ms | 0.27% | 111.686 ms | 0.35% | -3661.206 us | -3.17% | FAST |
| STRUCT | DEVICE_BUFFER | 0 | 32 | 500000 | 536870912 | 0 | 0 | 164.191 ms | 0.45% | 159.679 ms | 0.38% | -4511.743 us | -2.75% | FAST |
| STRUCT | DEVICE_BUFFER | 1000 | 32 | 500000 | 536870912 | 0 | 0 | 118.041 ms | 0.26% | 115.891 ms | 0.48% | -2150.305 us | -1.82% | FAST |
Summary
- Total Matches: 40
- Pass (diff <= min_noise): 1
- Unknown (infinite noise): 0
- Failure (diff > min_noise): 39
|
TPC-H is completely unaffected by these changes since there are no nontrivial repetition/definition levels to preprocess for those tables, which are all scalar types with no nulls. TPC-DS is slightly affected by this since there are some nullable columns in the tables, but there are still no nested data, so potential gains are limited. I did a bunch of runs of TPC-DS to try and measure something beyond just noise. I wasn't very scientific about it, and the results are technically still within the expected spread, so I hesitate to make any strong claims here, but averaged over 10 runs with 2 iterations each there is a small improvement of 1-2% that I don't think is just noise. The net effect remains quite small though. |
|
|
Ah good catch, sorry about that! I fixed it so the signs are consistent. |
| fill_index = 0; | ||
| decode_index = -1; // signals the first iteration. Nothing to decode. | ||
|
|
||
| cudf_assert(stage_capacity >= 0); |
There was a problem hiding this comment.
Ensure stage_capacity cannot exceed the buffer size for this specialization. Otherwise the async copy can write past the shared-memory buffer.
| cudf_assert(stage_capacity >= 0); | |
| cudf_assert(stage_capacity >= 0 and stage_capacity <= smem_stage_size); |
| // Slot 0 special case: resume a run that was split by the previous | ||
| // call. `cur` already points past this run's payload (fully | ||
| // consumed last call), so we do NOT re-parse its header - we just | ||
| // reuse the saved meta and continue emitting values. | ||
| if (partial_run_meta != -1) { | ||
| int const remaining = partial_run_total - partial_run_done; | ||
| int const room = out_end - out_base; | ||
| int const run_len = min(remaining, room); | ||
| chunk_meta_v[0] = partial_run_meta; | ||
| chunk_out_off_v[1] = run_len; | ||
| s_run0_payload_offset = partial_run_done; | ||
| num_runs = 1; | ||
| run_prefix_end = run_len; | ||
| if (run_len < remaining) { | ||
| partial_run_done += run_len; | ||
| } else { | ||
| partial_run_meta = -1; | ||
| } | ||
| } |
There was a problem hiding this comment.
AI pointed out that the partial-run resume path has no production coverage: preprocess_levels_kernel is the only chunked caller and invokes each decoder once. Add a focused test that splits both a long RLE run and a literal run across two decode_next calls, or explicitly restrict chunked decoding to a single call.
There was a problem hiding this comment.
I removed the partial run code for now. I may want it in a future PR building on top of this one, but I can add it back in then. We're not using it now and I'm not 100% sure it will prove useful later.
| while (written < output_count) { | ||
| int const batch_size = min(num_rle_stream_decode_threads, output_count - written); | ||
| if (t < batch_size) { | ||
| output[rolling_index<max_output_values>(cur_values + written + t)] = 0; |
There was a problem hiding this comment.
Please add regression coverage for the zero-bit fast path. decode_fixed.cu repeatedly decodes dictionary indices into a rolling buffer, and changing written + t to cur_values + written + t fixes stale slots after the first decode. Cover a nullable, single-value dictionary with more than rolling_buf_size values.
There was a problem hiding this comment.
I tried to add a test for this, but it turns out the level_bits == 0 fast path is never entered with cur_values > 0 in any reachable code path today. All calls stacks that reach here short-circuit on dict_bits == 0 before they reach this point. I verified this by injecting the bug back, removing the (dict_bits > 0) guard, and pre-poisoning sb->dict_idx[0..rolling_buf_size), and all PARQUET_TEST cases still pass. I just kept the defensive cur_values + written + t form for completeness.
| @@ -440,4 +748,7 @@ struct rle_stream { | |||
| __device__ inline int decode_next(int t) { return decode_next(t, max_output_values); } | |||
There was a problem hiding this comment.
decode_next_chunked does not update output_pos, but skip_runs relies on it. Let's add a static_assert to ensure that doesn't happen.
static_assert(not use_chunked_expand, "skip_decode is not supported by chunked-expand");| ©_barrier, | ||
| rle_stream_t::smem_stage_size); |
There was a problem hiding this comment.
Use the same rle_stream specialization for rle_stream_t as for decoders. The stage sizes match today, but a future template change could make this capacity disagree with the actual shared-memory buffer.
There was a problem hiding this comment.
Done in 44d0571dfd with the decoder_stream_t alias.
| __shared__ rle_run rep_runs[rle_run_buffer_size]; | ||
| static constexpr int max_output_values = cuda::std::numeric_limits<int>::max(); | ||
| rle_stream<level_t, level_decode_block_size, max_output_values> | ||
| rle_stream<level_t, level_decode_block_size, max_output_values, true> |
There was a problem hiding this comment.
Let's use the alias directly here instead of true
| rle_stream<level_t, level_decode_block_size, max_output_values, true> | |
| rle_stream_chunked<level_t, level_decode_block_size, max_output_values> |
There was a problem hiding this comment.
Done in 44d0571dfd with the decoder_stream_t alias introduced above.
| @@ -418,7 +418,7 @@ CUDF_KERNEL void __launch_bounds__(level_decode_block_size) | |||
| __shared__ rle_run def_runs[rle_run_buffer_size]; | |||
| __shared__ rle_run rep_runs[rle_run_buffer_size]; | |||
There was a problem hiding this comment.
decode_next_chunked does not use runs. Remove these shared ring buffers for chunked streams, or make the chunked stream constructible without an rle_run*, to recover shared memory.
There was a problem hiding this comment.
Done in ceb81d66f7. I went with the second option and split the constructor with requires clauses on use_chunked_expand.
| if (level_run & 1u) { | ||
| int const groups = level_run >> 1; | ||
| run_len = groups * 8; | ||
| run_desc = static_cast<int>(cur - s_start) | (1u << 31); |
There was a problem hiding this comment.
| run_desc = static_cast<int>(cur - s_start) | (1u << 31); | |
| run_desc = static_cast<uint32_t>(cur - s_start) | (1u << 31); |
| cur += groups * level_bits; | ||
| } else { | ||
| run_len = level_run >> 1; | ||
| run_desc = static_cast<int>(cur - s_start); |
There was a problem hiding this comment.
| run_desc = static_cast<int>(cur - s_start); | |
| run_desc = static_cast<uint32_t>(cur - s_start); |
| // Assumption: (cur - s_start) fits in 31 bits, i.e. the encoded | ||
| // level stream for a single Parquet page is < 2 GiB. This is | ||
| // guaranteed by the Parquet format in practice (page payloads are | ||
| // orders of magnitude smaller than 2 GiB) and is independent of |
There was a problem hiding this comment.
Optional nit: Perhaps we want to add a cudf_assert enforcing the 2GB limit?
mhaseeb123
added
3 - Ready for Review
The chunked-expand decoder had cross-call resume state (partial_run_meta / partial_run_total / partial_run_done plus s_run0_payload_offset) written in anticipation of a chunked-expand dict_stream in decode_page_data_generic that was later benchmarked and formally deferred (see the DEFER verdict recorded on feature branch opt/rle-def-rep-split, geomean B/A 1.0083x inside noise band with a Q18 regression). The current sole caller (preprocess_levels_kernel) invokes decode_next_chunked exactly once per stream with max_output_values = INT_MAX, so a single RLE run cannot exceed the output window and the resume path is structurally unreachable. Untested unreachable state-machine code was flagged in review; deleting it is the cleanest response. Replace the overflow-stash block with a cudf_assert enforcing the single-call invariant, and leave a comment referencing the prior chunked-dict work (commits 4dbde92dd1 / 3e349acb27) as a recovery pointer for whoever revisits multi-call chunked decoding.
The rle_stream::decode_next fast path for level_bits == 0 uses `cur_values + written + t` as the ring index rather than the simpler `written + t`. No current caller enters this fast path with cur_values > 0 (the writer floors dict_rle_bits >= 1 in chunk_dict.cu, and REPETITION/DEFINITION decoders in decode_preprocess.cu are single-call), so the simpler form would also pass all end-to-end tests. Documenting the invariant explicitly so it is preserved defensively for any future caller that iterates decode_next with level_bits == 0.
Introduce class-level constants `run_desc_literal_flag` (1u << 31) and `run_desc_offset_mask` (0x7fffffffu) on rle_stream and use them at all five sites that previously spelled the literal-run flag / offset mask inline. Also folds the standalone block comment about the 2 GiB invariant into the constants' docblock. Pure naming refactor -- no behavior change. Verified PARQUET_TEST still passes 481/481.
3 participants
Description
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:
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.
Checklist