apply alp patches in parallel in cuda

[onursatici]

We used to walk through the patches after decoding ALP then apply them one by one. Now we scatter them among workers and apply these in parallel

[github-actions]

Polar Signals Profiling Results

Latest Run

Status	Commit	Job	Attempt	Link
🟢 Done	af05b28		1	Explore Profiling Data

Previous Runs (4)

Status	Commit	Job	Attempt	Link
🟢 Done	2cc5a01		1	Explore Profiling Data
🟢 Done	cb6b225		1	Explore Profiling Data
🟢 Done	3381fe1		1	Explore Profiling Data
🟢 Done	56c4603		1	Explore Profiling Data

---

Powered by Polar Signals Cloud

[github-actions]

Benchmarks: Vortex queries

Verdict: No clear signal (low confidence)
Attributed Vortex impact: +5.8%
Engines: DataFusion No clear signal (+7.4%, low confidence) · DuckDB No clear signal (+4.1%, low confidence)
Vortex (geomean): 1.036x ➖
Parquet (geomean): 1.007x ➖
Shifts: Parquet (control) +0.7% · Median polish +4.3%

How to read Verdict and Engines

• Verdict: Overall PR-level signal after subtracting baseline drift estimated from Parquet control rows. It can be Likely improvement, Likely regression, or No clear signal.
• Engines: Per-engine attribution. DataFusion is compared against DataFusion/Parquet controls; DuckDB is compared against DuckDB/Parquet controls. This answers whether each engine improved or regressed independently.
• Confidence: Based on directional consistency, share of rows above the noise floor, and control-run noise.

---

datafusion / vortex-file-compressed (1.067x ➖, 0↑ 0↓)

name	PR af05b28 (ns)	base 3975895 (ns)	ratio (PR/base)
vortex_q00/datafusion:vortex-file-compressed	9958929	9303558	1.07
vortex_q01/datafusion:vortex-file-compressed	6503645	6119813	1.06

datafusion / parquet (0.993x ➖, 0↑ 0↓)

name	PR af05b28 (ns)	base 3975895 (ns)	ratio (PR/base)
vortex_q00/datafusion:parquet	20323096	20667041	0.98
vortex_q01/datafusion:parquet	4617335	4604281	1.00

duckdb / vortex-file-compressed (1.064x ➖, 0↑ 0↓)

name	PR af05b28 (ns)	base 3975895 (ns)	ratio (PR/base)
vortex_q00/duckdb:vortex-file-compressed	10585279	9980727	1.06
vortex_q01/duckdb:vortex-file-compressed	6357106	5950906	1.07

duckdb / parquet (1.022x ➖, 0↑ 0↓)

name	PR af05b28 (ns)	base 3975895 (ns)	ratio (PR/base)
vortex_q00/duckdb:parquet	23566612	23222387	1.01
vortex_q01/duckdb:parquet	9580946	9308681	1.03

---

No file size changes detected.

[0ax1]

@claude review this pr. and check whether scatter_alp_patches_range is even specific to alp

[joseph-isaacs]

force inline?

[0ax1]

yeah no force inline please, inline should always be enough

[github-actions]

Claude finished @0ax1's task in 3m 13s —— View job

---

PR Review: apply alp patches in parallel in cuda

• Read the changed CUDA kernel (dynamic_dispatch.cu)
• Read the benchmark changes (dynamic_dispatch_cuda.rs)
• Check whether scatter_alp_patches_range is actually specific to ALP
• Address the force inline? review comment

Is scatter_alp_patches_range specific to ALP?

No — the function body has nothing ALP-specific in it. It operates purely on generic FastLanes patch machinery: GPUPatches, PatchesCursor<T>, and FL_CHUNK. These are the same primitives used for bit-unpacking exception patches (BitunpackParams::patches_ptr, scatter_patches_chunk at line 309). There is no ALP exponent/float math anywhere inside it.

The only genuinely ALP-specific code is the caller, scatter_scalar_patches (line 244), which selects the payload with op.op_code == ScalarOp::ALP. That's the right place for the ALP knowledge to live.

So the alp in the name is a misnomer. Two concrete suggestions:

• Rename to scatter_patches_range to mirror the existing scatter_patches_chunk and reflect that it's encoding-agnostic. The doc comment ("Scatter ALP patches overlapping…") should drop "ALP" too. Fix this →
• Note that scatter_patches_chunk (line 194) is essentially a special case of the new range variant (single chunk, whole chunk, no bounds check). They can't be trivially merged because scatter_patches_chunk addresses out chunk-locally (out[patch.index]) while the range version addresses absolutely, but it's worth a follow-up to see if one can be expressed in terms of the other.

force inline? (joseph-isaacs, line 210)

Already handled — the function is declared __device__ __forceinline__, consistent with scatter_patches_chunk. So the answer is "yes, and it already is." No change needed unless you specifically want to stop force-inlining it (it does contain a nested loop, but a device-side scatter helper is exactly where you want no call overhead).

Correctness — looks sound

I traced the coordinate math against the old per-element path and it's equivalent:

• original_pos = original_chunk * FL_CHUNK + patch.index reconstructs the patch's absolute original position, then out[original_pos - patches.offset] maps back to the sliced/logical output index. This matches the old orig = my_pos + patches.offset / within = orig % FL_CHUNK logic exactly, including the sliced-offset case (patches.offset % FL_CHUNK != 0), since PatchesCursor is opened with the same chunk = original_chunk - patch_chunk_start.
• The original_pos >= original_start && original_pos < original_end bounds check is necessary and correct: a boundary FL chunk can hold patches belonging to an adjacent block, and this filters them so blocks never write outside their [block_start, block_start+block_len) range. No cross-block clobbering.
• Synchronization is correct. The two __syncthreads() around the scatter (line 247/250) fence the streaming __stcs writes (output stage) / in-place smem writes (input stage) before exceptions overwrite them, and re-fence so a later stage can consume patched smem. __syncthreads() carries a block-scope memory fence, so the ordinary write is visible before the patch overwrite even though a different thread may own each. The branch is uniform across the block (op.op_code/patches_ptr are the same for all threads), so the barriers aren't divergent.

Performance

This is the actual win: the old path had every value linearly scan its chunk's patch list looking for a match (O(values × patches_in_chunk)), whereas the new path divides each chunk's patches across all threads via PatchesCursor(patches, chunk, threadIdx.x, blockDim.x) and scatters them cooperatively. Exceptions are sparse, so this collapses the exception handling to roughly O(patches / threads) per chunk. Good change, and it aligns with the repo's "don't run a per-element accessor in an O(n) loop" guidance.

Minor: scatter_scalar_patches issues two __syncthreads() per ALP op. In practice there's at most one ALP op per stage so this is fine; if a stage could ever carry multiple patched ops, the paired barriers per op are slightly more than needed, but not incorrect.

Benchmark change

The dynamic_dispatch_cuda.rs change is a nice improvement: it now parameterizes bench_alp_for_bitpacked_f64 over (None, Some(100)) to add a 1%-patches variant (injecting PI every 100th value), which is exactly the case this PR optimizes — the no-patches variant would show no signal. Using into_parts() to carry the real ALP patches through into the reconstructed tree, plus the assert!(bp.patches().is_none(), "expected only ALP patches") guard, correctly isolates ALP exception patches from bit-packing patches so the benchmark measures the intended path.

Summary

Correct, well-synchronized, and a real algorithmic improvement to exception handling. The one substantive nit is the name: scatter_alp_patches_range (and its doc comment) are not ALP-specific and would be clearer as scatter_patches_range, matching scatter_patches_chunk. Everything else is optional polish. I did not run the CUDA build/tests (no GPU in this environment) — worth confirming vortex-cuda kernels compile and the benchmark builds locally.

[codspeed-hq]

Merging this PR will not alter performance

⚠️ Unknown Walltime execution environment detected

Using the Walltime instrument on standard Hosted Runners will lead to inconsistent data.

For the most accurate results, we recommend using CodSpeed Macro Runners: bare-metal machines fine-tuned for performance measurement consistency.

✅ 1712 untouched benchmarks
🆕 1 new benchmark
⏩ 10 skipped benchmarks¹

Performance Changes

	Mode	Benchmark	BASE	HEAD	Efficiency
🆕	WallTime	cuda/alp_for_bp_6bw_f64_1pct_patches/dispatch_f64[100M]	N/A	2.2 ms	N/A

---

_{Comparing os/alp-patches (af05b28) with develop (3975895)}

Footnotes

1. 10 benchmarks were skipped, so the baseline results were used instead. If they were deleted from the codebase, click here and archive them to remove them from the performance reports. ↩

[0ax1]

Great perf boost! Let's land this once we change the name for scatter_alp_patches_range and change the force inline to a simple inline.