Repository navigation
fix(cuda): bound dictionary gathers by the dictionary length - #10915
Conversation
Codes under null positions may hold any value, and so may corrupt input. The CUDA dictionary gathers read values[code] for every row, so such a code read outside the dictionary. Every gather now checks the code against the dictionary length and writes a zero value for a code outside it. Adds tests for out-of-range codes under nulls for string, primitive, decimal and bool values (with host-resident and device-resident codes), the dictionary the GPU fuzzer crashed on, all-null codes, codes sliced across 1024/2048 boundaries, full-range u16 and beyond-u16 dictionaries, the 12-byte inline boundary, values longer than 4 KiB and non-UTF-8 binary values. Co-Authored-By: Claude Opus 5.5 <[email protected]> Signed-off-by: Tanish Gudise <[email protected]>
The GPU fuzz target only generated dictionaries with primitive values. It now also generates dictionaries with UTF-8 or binary values, dictionaries whose nullable codes hold arbitrary values under null positions, and string arrays compressed with the CUDA preset. Co-Authored-By: Claude Opus 5.5 <[email protected]> Signed-off-by: Tanish Gudise <[email protected]>
|
@TanishGudise Thanks for having a stab at Vortex CUDA! I'll have a peek at your PRs in the next days. |
|
Thanks Alexander! Happy to make any changes, and I can rerun the GPU tests on my side if that helps. |
Merging this PR will degrade performance by 12.52%2 warnings
Warning Please fix the performance issues or acknowledge them on CodSpeed. Performance Changes
Tip Investigate this regression by commenting Comparing Footnotes
|
|
@TanishGudise Having a look just now |
0ax1
left a comment
There was a problem hiding this comment.
Thanks again for setting up the PR!
2 nits inline
| /// Dictionary encoding with primitive values. | ||
| Dict, | ||
| /// Dictionary encoding with UTF-8 or binary values. | ||
| StringDict, |
There was a problem hiding this comment.
I would suggest to split out all fuzz test logic unrelated to this PR. How about keeping DictGarbageNullCodes in this PR, and moving all logic related to StringDict and CudaCompressedStrings to a diff PR? This isn't really related to the fix.
There was a problem hiding this comment.
Done in 2bc042b. Kept DictGarbageNullCodes here and moved the StringDict/CudaCompressedStrings fuzz logic to a separate PR to follow.
| // checks the code against the dictionary length. | ||
| template <typename IndexT> | ||
| __device__ inline bool in_dictionary(IndexT code, uint64_t values_len) { | ||
| return static_cast<uint64_t>(code) < values_len; |
There was a problem hiding this comment.
I would suggest inlining this return static_cast<uint64_t>(code) < values_len;. It's not completely guessable what in_dictionary is.
There was a problem hiding this comment.
Done in 25b4f44, inlined at both call sites.
Replace the `in_dictionary` helper with the comparison it wrapped at its two call sites, so the check reads directly in each kernel. Co-Authored-By: Claude Opus 5.5 <[email protected]> Signed-off-by: Tanish Gudise <[email protected]>
The string-dictionary and CUDA-preset string generators in the GPU fuzz target aren't part of the bounds fix, so they move to a separate change. The fuzz target keeps the dictionaries with arbitrary codes under null positions, which exercise the bounds check, and the regression test now describes the dictionary the fuzzer found without quoting input bytes that no longer reproduce it. Co-Authored-By: Claude Opus 5.5 <[email protected]> Signed-off-by: Tanish Gudise <[email protected]>
0ax1
left a comment
There was a problem hiding this comment.
Thanks for addressing the comments, let's land this!
…eset (#10917) ## Summary - Tracking Issue: #10914 - Depends on #10915 and #10916; this branch contains both. The diff of its own is one line plus a test. `CompressionMode::Cuda` excludes `StringDictScheme`, but `execute_dict_varbinview` already decodes string dictionaries on the GPU. Keeping the scheme makes GPU-mode files 16.0% smaller across compress-bench's 17-dataset GPU suite (3,351.3 → 2,813.9 MB). Per-dataset gains: CMSprovider 45.7%, Food 52.9%, HashTags 33.2%, both TPC-H l_comment variants 12.2%, Euro2016 10.7%. The issue has the full table. ## Changes - `vortex-btrblocks/src/builder.rs`: drops `StringDictScheme` from the Cuda exclusions and updates the comment. A new test checks that the preset keeps string dictionaries but not binary ones. This is the only change to CPU-compressor code. Evidence with this branch (A40, CUDA 13.0): - `--gpu-verify` passes on all 17 datasets. develop plus only this line fails 5: taxi, CMSprovider, Euro2016, Food, HashTags. - The CPU-fallback census is identical to develop's: 59 constant-array fallbacks, none from dictionaries. - The dictionary gather takes at most 6.1% of decode time and reaches 62-85% of the measured 563 GB/s. - `vortex-cuda` lib suite on this branch: 723 passed, 0 failed. - `compress_gpu` fuzz (CI flags, ASAN): 100,000 runs in 256 s, 0 crashes. ## Benchmarks Setup: - `compress-bench --gpu-decompress --iterations 5`, a fresh process per run, `taskset -c 2-7`, 3 ABBA rounds. - Speedup = develop time / branch time. - The same-design A/A floor is 0.14 in |log|, so a "faster" claim needs a 95% CI above 1 and a speedup above 1.32×. | Dataset | File MB, develop → branch | Speedup [95% CI] | | |---|---|---:|---| | Food | 84.7 → 39.9 | 1.55× [1.41, 1.71] | faster | | CMSprovider | 771.0 → 418.5 | 1.37× [0.66, 2.86] | within noise | | Euro2016 | 163.3 → 145.9 | 1.46× [0.92, 2.32] | within noise | | TPC-H l_comment chunked | 169.6 → 148.8 | 1.36× [0.55, 3.34] | within noise | | taxi | 52.0 → 51.5 | 1.28× [0.32, 5.11] | within noise | | HashTags | 235.2 → 157.0 | 1.22× [0.89, 1.69] | within noise | | TPC-H l_comment canonical | 169.6 → 148.8 | 1.18× [0.71, 1.96] | within noise | | Arade | 144.2 → 141.7 | 1.05× [0.42, 2.68] | within noise | The nine datasets whose files don't change got one AB/BA round each. Eight came out at 1.01-1.18×, and one (wide table cols=1000 chunks=50) at 0.91×, inside the A/A floor. The measurements are from an A40 (GA102); the repo's GPU benchmarks run on A10G. ## API changes `only_cuda_compatible()`, and therefore `cuda_write_strategy`, can now emit string dictionaries. - An `ArrowDeviceArrayStream` with the default `DictionaryExport::Preserve` rejects a batch whose Arrow schema differs from the first. - With this change, 6 of CMSprovider's 14 string columns mix dictionary and constant chunks. - Today's Cuda files already do this in numeric columns: 7 of taxi's 18 numeric columns and 5 of CMSprovider's 12. - Streams with `DictionaryExport::Decode`, which the FFI scan uses, are unaffected. Overlap: #10425 adds a line next to the removed one, so whichever merges second needs a one-line rebase. #9762 changes dictionary selection, so the sizes are as of a59294e. ## AI assistance Developed with Claude Code (Claude Opus 5.5), which wrote most of the code, tests and this description under my direction. I reviewed the diff and the measurements. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Signed-off-by: Tanish Gudise <[email protected]> Co-authored-by: Claude Opus 5.5 <[email protected]> Co-authored-by: Alexander Droste <[email protected]>
Summary
The standalone dictionary kernels in
dict.cugatheredvalues[codes[i]]for every row. Nullable codes can hold any value at null positions (DictArray::try_newdoesn't bound-check them), so these gathers read outside the dictionary. The fused planner already rejects nullable-code dictionaries for this reason and sends them to these kernels. Develop's own fuzzer, extended here, crashes on this within 71k runs.Changes
dict.cu: every kernel (primitive, decimal, bool and the 16-byte view gather) takesvalues_len. A code at or beyond it writes a zero value. In a valid array only null positions can hold such a code, so the zero is never observed.dict.rsandbenches/dynamic_dispatch_cuda.rs: passvalues_lenat every launch.assert_arrays_eq:u32::MAXandu64::MAX;fuzz/src/gpu/mod.rs:compress_gpualso generates string and binary dictionaries, dictionaries with arbitrary u8/u16/u32 codes under nulls, and strings compressed with the CUDA preset.Evidence (A40, CUDA 13.0):
CUDA_ERROR_ILLEGAL_ADDRESScompute-sanitizer --tool memcheckon the 6 out-of-range testsInvalid __global__ readof 1, 8 or 16 bytescompress_gpufuzz (CI flags, ASAN, empty corpus)CUDA_ERROR_ILLEGAL_ADDRESSvortex-cudalib suiteUnder nsys, the bounds check costs nothing measurable on numeric dictionaries: 1.574 vs 1.563 ms per iteration on CMSprovider and 0.825 vs 0.832 ms on taxi.
On develop's kernels the
u32::MAX/u64::MAXtests leave the CUDA context in a stickyCUDA_ERROR_ILLEGAL_ADDRESSstate. Run in one process, they also fail the unrelated tests that follow, about 336 of them. A regression of this fix would show up the same way, as a wall of failures led by these tests.Everything in the table ran on this PR alone, except the develop fuzz crash: that ran develop's kernels under this PR's extended fuzz target.
AI assistance
Developed with Claude Code (Claude Opus 5.5), which wrote most of the code, tests and this description under my direction. I reviewed the diff and the measurements.
🤖 Generated with Claude Code