CUDA: two GPU decode bugs on develop, and string dictionaries in the CUDA preset (16% smaller GPU files) #10914

Two GPU decode bugs reproduce on develop @ a59294e. Both also block allowing StringDictScheme in CompressionMode::Cuda, which would make GPU-mode files 16.0% smaller on compress-bench's GPU suite. I have three small PRs ready and one question about #8921 at the end.

Bug 1: into_host panics on validity that CUDA execution left encoded

cargo run -p compress-bench --profile release_debug --features cuda -- --datasets '^CMSprovider$' --gpu-decompress --iterations 1 --gpu-verify

GPU mode also runs cuDF for the Parquet side, so the Python on PATH needs cudf-cu12 (I used 26.8). On an A40 with CUDA 13.0, this panics with expected host buffer, and Euro2016 panics the same way:

vortex_cuda::canonical::validity_into_host
  -> vortex_cuda::canonical::child_into_host
  -> <vortex_runend::array::RunEnd as VTable>::execute
  -> vortex_runend::decompress_bool::runend_decode_bools
  -> BufferHandle::as_host

A run-end encoded validity mask stays encoded on the device after execute_cuda. CanonicalCudaExt::into_host then executes it on the host while its buffers are still on the device.

A nullable dictionary does the same thing. The four GPU dict executors return the output validity as values_validity.take(codes), a lazy array over device-resident codes. The existing dict tests pass host-resident codes, so they don't reach this path.

Bug 2: dictionary gathers read out of bounds at null positions

The standalone kernels in vortex-cuda/kernels/src/dict.cu gather values[codes[i]] for every row. DictArray::try_new doesn't bound-check codes, and nullable codes can hold any value at null positions. The fused planner already guards against this: plan_builder.rs rejects a Dict with nullable codes because "Nullable codes could hold garbage values at null positions, causing out-of-bounds shared memory reads". The array then goes to the standalone kernel, which has no guard.

On develop's kernels:

• compute-sanitizer --tool memcheck reports Invalid __global__ read of 1, 8 or 16 bytes, 1-2 errors per test, on string, decimal, bool and primitive dictionaries with out-of-range codes under nulls.
• Codes like u32::MAX under nulls fail outright with CUDA_ERROR_ILLEGAL_ADDRESS.
• The compress_gpu fuzz target, extended to generate garbage null codes, crashes after 70,941 runs with CUDA_ERROR_ILLEGAL_ADDRESS (CI flags, ASAN). The input is a one-value u8 dictionary whose null code is 3,960,089,894.

What the bugs block: string dictionaries in the CUDA preset

only_cuda_compatible() excludes StringDictScheme, although execute_dict_varbinview (#6361) already decodes string dictionaries on the GPU. With the scheme allowed and both bugs fixed, the 17 datasets of compress-bench's GPU suite come out as follows:

Sizes are exact file bytes at a59294e; the last column uses the PR 3 preset. String dictionaries recover 98% of the size the Cuda preset costs relative to Default.

Results with the change:

• --gpu-verify passes on all 17 datasets.
• The RUST_LOG=vortex_cuda=debug fallback census is identical to develop's: 59 constant-array fallbacks, none from dictionaries.
• Decode is no slower. Food is faster, 1.55× [1.41, 1.71]; the other changed datasets are within noise.

Plan: three PRs

1. fix(cuda): bound dictionary gathers by the dictionary length #10915 fix(cuda): bound dictionary gathers by the dictionary length.
   • Every dict.cu kernel takes values_len, and a code outside the dictionary writes zero.
   • Tests cover out-of-range codes under nulls, including the fuzzer's input.
   • The compress_gpu target is extended to generate string dictionaries and garbage null codes.
2. fix(cuda): copy encoded device-resident children to the host in into_host #10916 fix(cuda): copy encoded device-resident children to the host in into_host.
   • An encoded child's device buffers are copied to the host before it is executed there.
   • Tests cover run-end and sparse masks, and string and primitive dictionaries with nullable values and device codes.
   • gpu-scan-cli scan --gpu-file panics with "as_slice must be called on host buffer" #7838 says mixed CPU/GPU execution isn't supported yet. This doesn't add it: the whole child moves to the host first, so nothing decodes from device buffers on the host.
3. feat(btrblocks): allow string dictionaries in the CUDA compression preset #10917 feat(btrblocks): allow string dictionaries in the CUDA preset.
   • The one exclusion line, stacked on 1 and 2.

I also tried a different fix for bug 1: materializing validity in the dict executors with execute_validity_cuda. It doesn't cover run-end masks. For a dictionary with both nullable codes and nullable values, it errors, because vortex.fill_null has no CUDA kernel.

Also found, not fixed:

• Signed dictionary codes have no GPU kernels: dict.cu instantiates u8/u16/u32/u64 only, so i8 codes fail with Failed to load kernel function 'dict_f32_i8'.
• vortex.fill_null has no CUDA kernel.

Question

@0ax1, #8921 added the preset comment "Other string fragmentation and dictionary schemes still require unsupported decode paths". Which path did you mean for string dictionaries?

Beyond the two bugs, nothing else failed in these checks:

• --gpu-verify on 17 datasets;
• the fallback census;
• 16 new tests and the full vortex-cuda lib suite, on each PR alone;
• 100,000 fuzz runs with 0 crashes.

My guess is the Arrow device stream.

• With the default DictionaryExport::Preserve, ArrowDeviceArrayStream rejects a batch whose schema changes ("chunks must not vary their encoding"). test_default_dictionary_device_stream pins that behaviour.
• With StringDict allowed, 6 of CMSprovider's 14 string columns mix dictionary and constant chunks.
• Today's preset already does this for numeric columns: 7 of taxi's 18 numeric columns and 5 of CMSprovider's 12 mix dictionary chunks with FoR, BitPacked or ALP chunks.
• DictionaryExport::Decode, which the FFI scan uses, is unaffected.

PR 3 documents this as an API change rather than changing the stream. Is that the path you meant, or is there another one?