Skip to content

Optimize RLE decoding using a warp-balanced chunking approach #23271

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

Open
vyasr wants to merge 37 commits into
base: main
Choose a base branch
from
+362 −48
Open

Optimize RLE decoding using a warp-balanced chunking approach #23271

Show file tree
Hide file tree
Changes from 26 commits
Commits
Show all changes
37 commits
Select commit Hold shift + click to select a range
5fec7c1
perf(parquet): add chunked-parse expand path to rle_stream (opt-in)
vyasr Jul 10, 2026
0a9f45a
perf(parquet): use chunked-parse expand in preprocess_levels_kernel
vyasr Jul 10, 2026
3f3e4fc
docs(parquet): document decode_fixed SMEM constraint on chunked RLE
vyasr Jul 10, 2026
910da6e
fix(parquet): fix signed UB in rle_stream chunked meta flag; add boun…
vyasr Jul 10, 2026
401cdfc
refactor(parquet): use cuda::std::span and upper_bound in rle_stream …
vyasr Jul 14, 2026
bb226c8
refactor(parquet): remove unused chunked_smem_t struct
vyasr Jul 14, 2026
c60de71
Some minor cleanup
vyasr Jul 14, 2026
c3f6c05
More comments
vyasr Jul 14, 2026
1e5d130
Hoist level_bits==0 fast path into decode_next
vyasr Jul 14, 2026
8f2b7f4
test(parquet): pass cooperative group to rle_stream::init in equivale…
vyasr Jul 14, 2026
042d8c5
perf(parquet): raise kGenRuns to 1024 on sm_80+ for chunked RLE expand
vyasr Jul 14, 2026
be9a068
docs(parquet): note the 2 GiB payload-offset assumption in gen_meta
vyasr Jul 14, 2026
c4facac
refactor(parquet): hoist level_mask out of chunked-expand literal loop
vyasr Jul 14, 2026
fca7320
perf(parquet): use fixed kGenRuns=512 across all archs
vyasr Jul 15, 2026
3388901
style(parquet): use snake_case for new chunked-expand identifiers
vyasr Jul 15, 2026
e2c899a
refactor(parquet): inline warp_fill into its sole caller
vyasr Jul 15, 2026
5396008
perf(parquet): use max_runs_per_chunk=1024 on sm_80+
vyasr Jul 15, 2026
2ce885c
refactor(parquet): template-parameterize rle_stream stage sizes and s…
vyasr Jul 17, 2026
317b9ac
Merge remote-tracking branch 'upstream/main' into opt/rle-chunked-expand
vyasr Jul 17, 2026
4317bd2
Merge branch 'main' into opt/rle-chunked-expand
mhaseeb123 Jul 23, 2026
ba909e8
Address PR review nits in RLE chunked decode
vyasr Jul 25, 2026
85b9532
Remove RLE chunked-expand equivalence test
vyasr Jul 25, 2026
041c4ab
Use cooperative groups in decode_next_chunked
vyasr Jul 25, 2026
18b7d0b
Use cooperative groups in decode_next_ring
vyasr Jul 25, 2026
798d545
Per-lane linear-walk phase-2 expand in decode_next_chunked
vyasr Jul 26, 2026
9e9d8bf
Use stable owners for RLE stream state
vyasr Jul 27, 2026
b120823
Merge branch 'main' into opt/rle-chunked-expand
mhaseeb123 Jul 29, 2026
b3515c7
Merge remote-tracking branch 'upstream/main' into opt/rle-chunked-expand
vyasr Jul 30, 2026
64c6af2
Address safety and correctness in edge cases raised during PR review
vyasr Jul 31, 2026
e3738f3
Fix signedness issues with run_desc
vyasr Jul 31, 2026
aebcc26
Various style and code quality improvements from PR feedback
vyasr Jul 31, 2026
44d0571
Single-source rle_stream type alias
vyasr Jul 31, 2026
ceb81d6
Don't waste runs shmem on chunked expansion rle_streams
vyasr Jul 31, 2026
789b511
Remove untested cross-call partial-run resume from chunked rle_stream
vyasr Jul 31, 2026
9128028
docs(parquet): explain level_bits==0 fast path cur_values offset
vyasr Jul 31, 2026
27446f9
refactor(parquet): name run_desc bit-packing constants
vyasr Aug 2, 2026
9a55599
Merge remote-tracking branch 'upstream/main' into opt/rle-chunked-expand
vyasr Aug 2, 2026
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
8 changes: 5 additions & 3 deletions cpp/src/io/parquet/decode_preprocess.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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];

@mhaseeb123 mhaseeb123 Jul 29, 2026

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

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.

@vyasr vyasr Aug 2, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Done in ceb81d66f7. I went with the second option and split the constructor with requires clauses on use_chunked_expand.

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>

@mhaseeb123 mhaseeb123 Jul 29, 2026

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

Let's use the alias directly here instead of true

Suggested change
rle_stream<level_t, level_decode_block_size, max_output_values, true>
rle_stream_chunked<level_t, level_decode_block_size, max_output_values>

@vyasr vyasr Aug 2, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Done in 44d0571dfd with the decoder_stream_t alias introduced above.

decoders[level_type::NUM_LEVEL_TYPES] = {{def_runs}, {rep_runs}};

// Shared-memory staging scratch for the encoded level streams. Level streams
Expand Down Expand Up @@ -452,7 +452,8 @@ CUDF_KERNEL void __launch_bounds__(level_decode_block_size)
rep,
num_to_decode,
stage,
&copy_barrier);
&copy_barrier,
rle_stream_t::smem_stage_size);

@mhaseeb123 mhaseeb123 Jul 28, 2026

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

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.

@vyasr vyasr Aug 2, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Done in 44d0571dfd with the decoder_stream_t alias.

copy_barrier.arrive_and_wait();
decoders[level_type::REPETITION].decode_next(t, num_to_decode);
}
Expand All @@ -475,7 +476,8 @@ CUDF_KERNEL void __launch_bounds__(level_decode_block_size)
def,
num_to_decode,
stage,
&copy_barrier);
&copy_barrier,
rle_stream_t::smem_stage_size);
copy_barrier.arrive_and_wait();
decoders[level_type::DEFINITION].decode_next(t, num_to_decode);
}
Expand Down
Loading
Loading