Skip to content

fix(cuda): bound dictionary gathers by the dictionary length - #10915

Merged
0ax1 merged 5 commits into
vortex-data:developfrom
TanishGudise:tg/pr1-dict-bounds
Nov 10, 2026
Merged

0ax1 merged 5 commits into
vortex-data:developfrom
TanishGudise:tg/pr1-dict-bounds

Conversation

@TanishGudise

Copy link
Copy Markdown
Contributor

Summary

The standalone dictionary kernels in dict.cu gathered values[codes[i]] for every row. Nullable codes can hold any value at null positions (DictArray::try_new doesn'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) takes values_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.rs and benches/dynamic_dispatch_cuda.rs: pass values_len at every launch.
  • Tests (12), each compared with the expected result through assert_arrays_eq:
    • out-of-range codes under nulls for string, decimal, bool and primitive values, with codes up to u32::MAX and u64::MAX;
    • the dictionary the fuzzer crashed on (one u8 value, null code 3,960,089,894);
    • device-resident codes with nullable values and codes;
    • all-null codes;
    • slices across 1,023/1,024/1,025 and 2,047/2,048;
    • u16 codes with 65,536 values and u32 codes with 70,000;
    • empty, 12-, 13- and 4,097-byte strings;
    • non-UTF-8 binary values.
  • fuzz/src/gpu/mod.rs: compress_gpu also 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):

Check develop's kernels this PR
5 bug tests: out-of-range codes under nulls (string, decimal, bool, primitive) and the fuzzer's input 5 fail with CUDA_ERROR_ILLEGAL_ADDRESS 5 pass
7 edge-case tests (all-null codes, block-boundary slices, full u16/u32 ranges, inline boundary, binary, device codes) pass 7 pass
compute-sanitizer --tool memcheck on the 6 out-of-range tests 1-2 errors each, Invalid __global__ read of 1, 8 or 16 bytes 0 errors
compress_gpu fuzz (CI flags, ASAN, empty corpus) crash after 70,941 runs, CUDA_ERROR_ILLEGAL_ADDRESS 100,000 runs in 283 s, 0 crashes
vortex-cuda lib suite — 719 passed, 0 failed

Under 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::MAX tests leave the CUDA context in a sticky CUDA_ERROR_ILLEGAL_ADDRESS state. 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

TanishGudise and others added 2 commits October 27, 2026 13:50
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]>
@0ax1

0ax1 commented Oct 28, 2026

Copy link
Copy Markdown
Contributor

@TanishGudise Thanks for having a stab at Vortex CUDA! I'll have a peek at your PRs in the next days.

@TanishGudise

Copy link
Copy Markdown
Contributor Author

Thanks Alexander! Happy to make any changes, and I can rerun the GPU tests on my side if that helps.

@codspeed

codspeed Bot commented Oct 28, 2026 •

Copy link
Copy Markdown

Merging this PR will degrade performance by 12.52%

2 warnings

⚠️ Walltime results may be unreliable

Some benchmarks used the Walltime instrument outside CodSpeed Macro Runners, which can make the results inconsistent.

Learn more about Macro Runners

See affected benchmarks

⚠️ 1 benchmark is dominated by system calls

System calls cannot be simulated consistently and are left out of the measure. The reported value understates the real cost. Switch to the Walltime instrument to measure them.

Learn more about system calls

See affected benchmarks

❌ 1 regressed benchmark
✅ 1962 untouched benchmarks
⏩ 1 skipped benchmark1

Warning

Please fix the performance issues or acknowledge them on CodSpeed.

Performance Changes

Mode Benchmark BASE HEAD Efficiency
❌ Walltime lt_i64_nullable_avx2 8.3 µs 9.5 µs -12.52%

Tip

Investigate this regression by commenting @codspeedbot fix this regression on this PR, or directly use the CodSpeed MCP with your agent.


Comparing TanishGudise:tg/pr1-dict-bounds (2bc042b) with develop (7b36f7f)

Open in CodSpeed

Footnotes

  1. 1 benchmark was skipped, so the baseline result was used instead. If it was deleted from the codebase, click here and archive it to remove it from the performance reports. ↩

@joseph-isaacs
joseph-isaacs requested a review from 0ax1 October 29, 2026 15:42
@0ax1 0ax1 added the changelog/fix A bug fix label Nov 6, 2026
@0ax1

0ax1 commented Nov 9, 2026

Copy link
Copy Markdown
Contributor

@TanishGudise Having a look just now

@0ax1 0ax1 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Thanks again for setting up the PR!

2 nits inline

Comment thread fuzz/src/gpu/mod.rs Outdated
/// Dictionary encoding with primitive values.
Dict,
/// Dictionary encoding with UTF-8 or binary values.
StringDict,

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

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.

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 2bc042b. Kept DictGarbageNullCodes here and moved the StringDict/CudaCompressedStrings fuzz logic to a separate PR to follow.

Comment thread vortex-cuda/kernels/src/dict.cu Outdated
// 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;

@0ax1 0ax1 Nov 9, 2026 •

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

I would suggest inlining this return static_cast<uint64_t>(code) < values_len;. It's not completely guessable what in_dictionary is.

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 25b4f44, inlined at both call sites.

TanishGudise and others added 2 commits November 9, 2026 16:43
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 0ax1 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Thanks for addressing the comments, let's land this!

@0ax1
0ax1 enabled auto-merge (squash) November 10, 2026 11:17
@0ax1
0ax1 merged commit 30446ea into vortex-data:develop Nov 10, 2026
79 of 80 checks passed
0ax1 added a commit that referenced this pull request Nov 10, 2026
…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]>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

changelog/fix A bug fix

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants