Skip to content

Feature/rle pipeline smem staging#8

Open
vyasr wants to merge 6 commits into
opt/rle-chunked-expandfrom
feature/rle-pipeline-smem-staging
Open

Feature/rle pipeline smem staging#8
vyasr wants to merge 6 commits into
opt/rle-chunked-expandfrom
feature/rle-pipeline-smem-staging

Conversation

@vyasr

@vyasr vyasr commented Jul 18, 2026

Copy link
Copy Markdown
Owner

This PR implements a two-staged pipelining approach for filling shared memory buffers with run data for warps to consume. I did some coarse sweeps of how much smem to allocate to different stages, but the deltas were not huge, so there's probably limited room for this approach to help us without bigger kernel-level changes.

vyasr and others added 6 commits July 17, 2026 23:36
…next_chunked

Implements task D of the chunked-smem-staging plan: double-buffered SMEM staging via cuda::pipeline in rle_stream::decode_next_chunked, addressed by chunk_size (4096 bytes). Broadcasts thread-0's cur offset through __shared__ before collective pipeline ops to avoid per-thread cur divergence and pipeline deadlock.

Ultraworked with [Sisyphus](https://github.com/code-yeongyu/oh-my-openagent)

Co-authored-by: Sisyphus <clio-agent@sisyphuslabs.ai>
…_levels_kernel

Implements task E of the chunked-smem-staging plan: preprocess_levels_kernel allocates pipeline_shared_state<2> and wires it to rle_stream::set_pipeline for both REPETITION and DEFINITION decoders.

Ultraworked with [Sisyphus](https://github.com/code-yeongyu/oh-my-openagent)

Co-authored-by: Sisyphus <clio-agent@sisyphuslabs.ai>
Adds 4 targeted correctness tests for the chunked-RLE-staging path in preprocess_levels_kernel (Tasks D+E+F):
- ExactBoundary: page rows == N * chunk_size
- SingleLongLiteralRun: >32K rows in one RLE literal run
- SkipAhead: skip_rows before decode_next_chunked
- ZeroLengthPage: 0 rows after row-group filter
Ultraworked with [Sisyphus](https://github.com/code-yeongyu/oh-my-openagent)

Co-authored-by: Sisyphus <clio-agent@sisyphuslabs.ai>
…tream

Ultraworked with [Sisyphus](https://github.com/code-yeongyu/oh-my-openagent)

Co-authored-by: Sisyphus <clio-agent@sisyphuslabs.ai>
Split get_vlq32 into an offset+accessor primary and a thin pointer-cursor
adapter. Callers that read from a moving global pointer are unchanged;
the accessor form lets the chunked-expand path dispatch each byte
between the double-buffered SMEM staging chunks and the raw global-mem
fallback without duplicating the VLQ state machine.

Removes the get_vlq32_at lambda inside rle_stream_chunked::decode_next
in favor of the shared implementation.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant