feat(cuda): export FSST directly to Arrow varbin

[0ax1]

Summary

Decode CUDA FSST arrays directly into Arrow’s standard offset-based Utf8/Binary layout: i32 offsets plus one contiguous device values buffer.

Previously, FSST decoded to Utf8View/BinaryView. Since cuDF does not consume view arrays, it then had to materialize the same offset-based representation itself. Direct varbin export avoids that additional conversion and enables BenchPress GPU reads for Vortex string columns.

Behavior changes

• CUDA sessions now default variable-length exports to standard Arrow Utf8/Binary.

• Consumers that support views can opt in with:

  CudaSession::with_varbin_export_layout(VarBinExportLayout::VarBinView)

[github-actions]

Benchmarks: Vortex queries

Verdict: No clear signal (low confidence)
Attributed Vortex impact: +3.1%
Engines: DataFusion No clear signal (+5.1%, low confidence) · DuckDB No clear signal (+1.1%, low confidence)
Vortex (geomean): 1.017x ➖
Parquet (geomean): 1.002x ➖
Shifts: Parquet (control) +0.2% · Median polish +0.7%

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.056x ➖, 0↑ 0↓)

name	PR 2cc117e (ns)	base 54b045b (ns)	ratio (PR/base)
vortex_q00/datafusion:vortex-file-compressed	9927213	9302721	1.07
vortex_q01/datafusion:vortex-file-compressed	6331074	6055120	1.05

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

name	PR 2cc117e (ns)	base 54b045b (ns)	ratio (PR/base)
vortex_q00/datafusion:parquet	20178473	20194975	1.00
vortex_q01/datafusion:parquet	4551961	4500660	1.01

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

name	PR 2cc117e (ns)	base 54b045b (ns)	ratio (PR/base)
vortex_q00/duckdb:vortex-file-compressed	9956945	9926207	1.00
vortex_q01/duckdb:vortex-file-compressed	6006185	5918276	1.01

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

name	PR 2cc117e (ns)	base 54b045b (ns)	ratio (PR/base)
vortex_q00/duckdb:parquet	23287163	23264613	1.00
vortex_q01/duckdb:parquet	9351080	9395799	1.00

---

No file size changes detected.

[codspeed-hq]

Merging this PR will improve performance by 51.12%

⚠️ 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.

⚡ 6 improved benchmarks
✅ 1691 untouched benchmarks
🆕 5 new benchmarks
⏩ 10 skipped benchmarks¹

Performance Changes

	Mode	Benchmark	BASE	HEAD	Efficiency
⚡	WallTime	cuda/bitpacked_u8/unpack/3bw[100M]	646.3 µs	300.3 µs	×2.2
⚡	WallTime	cuda/bitpacked_patched_u8/unpack/1%[100M]	768.9 µs	381.8 µs	×2
⚡	WallTime	cuda/bitpacked_patched_u8/unpack/10%[100M]	828.6 µs	477.3 µs	+73.58%
⚡	WallTime	cuda/bitpacked_patched_u16/unpack/1%[100M]	712.3 µs	595.4 µs	+19.64%
⚡	WallTime	cuda/bitpacked_patched_u16/unpack/10%[100M]	831 µs	714.1 µs	+16.36%
⚡	WallTime	cuda/bitpacked_u16/unpack/5bw[100M]	634.8 µs	558.1 µs	+13.73%
🆕	WallTime	cuda/fsst/decompress_to_varbin[10M]	N/A	19.3 ms	N/A
🆕	WallTime	cuda/fsst/decompress_to_varbinview[10M]	N/A	19.9 ms	N/A
🆕	WallTime	cuda/fsst/export_binary_view[10M]	N/A	20.1 ms	N/A
🆕	WallTime	cuda/fsst/export_binary[10M]	N/A	19.4 ms	N/A
🆕	WallTime	cuda/fsst/export_utf8_view[10M]	N/A	20 ms	N/A

Tip

Curious why this is faster? Comment @codspeedbot explain why this is faster on this PR, or directly use the CodSpeed MCP with your agent.

---

_{Comparing ad/cuda-fsst-varbin-export (2cc117e) with develop (54b045b)}

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]

@claude review

[github-actions]

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

---

Claude PR Review

Reviewed the full diff against origin/develop (PR base 159207e). This is a well-structured change with strong test coverage — the FSST→Arrow varbin path, the in-kernel view construction, and the BitBuffer::sliced Miri fix all hang together coherently. Findings below, ordered by importance.

🔴 Correctness / behavior

1. PR description is stale — no per-array fallback or fsst_varbin_offsets_fit predicate exists. The description says "FSST arrays whose decoded size exceeds i32::MAX bytes fall back to the view layout per-array. Schema derivation and data export share one predicate (fsst_varbin_offsets_fit)." Neither is in the code:

   • grep fsst_varbin_offsets_fit → no matches.
   • The actual behavior (confirmed by test_oversized_fsst_varbin_export_errors) is that oversized exports hard-error with "FSST decoded size exceeds Arrow i32 offset range" rather than falling back. Commit 3db845d ("fall back to view layout…") appears to have been superseded. Please update the PR description so reviewers/consumers aren't misled about the fallback semantics.

2. New default (Utf8) makes previously-exportable large string columns fail. With VarBinExportLayout::VarBin now the default, any Utf8 array whose decoded heap exceeds i32::MAX bytes — which previously exported fine as Utf8View — now returns an error instead of a valid array. This is inherent to choosing i32-offset Utf8, but it's a sharper edge than "the default layout changed": it's a silent-until-runtime failure for >2 GB string batches. Two things worth confirming:

   • Is erroring (vs. the originally-planned per-array view fallback) the intended final behavior? If so, the "Behavior changes" section should call this out explicitly so downstream consumers know to opt into with_varbin_export_layout(VarBinView) for large columns.
   • arrow_device_export_field unconditionally returns DataType::Utf8 for the schema even when the data export will error. Schema-succeeds / data-errors is a defensible "stable schema" choice (and is what test_oversized_fsst_varbin_export_errors asserts), but it directly contradicts the description's "always resolves to the same layout" claim.

🟡 Kernel / soundness

3. fsst_varbin values buffer relies on uncompressed_lengths being exact. output_offsets is the host prefix-sum of uncompressed_lengths (nulls contribute 0), and total_size is the last offset. The kernel writes each non-null string into [offset[sid], offset[sid+1]). If a decoded string were ever shorter than its claimed length, the gap bytes stay uninitialized device memory yet still fall inside total_size and get exported. This matches the pre-existing view-kernel contract, so it's not a regression — just flagging the invariant since the offset-based buffer has no per-view length to mask it.

4. In-kernel view construction (fsst_write_view) looks correct. Each thread reads back only the bytes it just wrote (no cross-thread reads, write-before-read ordered within the thread). The outlined branch's (uint32_t)start truncation is safe because views are only built when total_size <= MAX_BUFFER_LEN on the Rust side, keeping start within u32. The int32_t → uint64_t read of output_offsets in fsst_decode_string is safe given offsets are validated non-negative on the host. Good.

🟢 BitBuffer::sliced Miri fix

5. The rewrite to iter_padded().take(n_words) → BufferMut::<u64> → truncate is the right approach, and bundling it here makes sense since decode_fsst_varbin_typed now feeds validity…execute_mask(...).into_bit_buffer().sliced() straight to device FFI, which requires bit-0 alignment. Test coverage (aligned / unaligned / empty / crosses_word_boundary) is exactly the right matrix. Nice.

⚪ Minor / nits

• decode_fsst_varbin computes dtype, validity, and len, then decode_fsst_varbin_typed re-derives all three from fsst again. Cheap (clones/getters), but the duplication could be threaded through instead.
• Offset prefix-sum is done on the host (O(n) per export). cub::exclusive_sum_i32 is already imported in canonical.rs — a device sum would keep the whole path GPU-resident, though it's not required and matches the existing view path.
• The assert_eq!(output_base_ptr % 16, 0, …) alignment check panics rather than returning VortexResult. Consistent with the existing decode_fsst kernel, so fine, but it is a panic in library code.

📊 CI signals

• codspeed rank_single[(1024, 0.9)] −10.96%: vortex-mask rank does not call BitBuffer::sliced (it uses iter_padded directly in intersect_by_rank.rs), and this PR touches no mask code. Combined with codspeed's own "Unknown Walltime execution environment / Hosted Runners lead to inconsistent data" warning, this reads as measurement noise rather than a real regression. Worth a glance but likely ignorable.
• "SQL BENCHMARK FAILED": please confirm this is infra/unrelated and not triggered by the default Utf8View→Utf8 export change (a consumer expecting the view layout could break).

Overall: solid, thoroughly-tested change. The main asks are (1) fix the stale PR description and (2) confirm the oversized-export error (vs. fallback) is the intended contract and document the large-column implication.
· ad/cuda-fsst-varbin-export

[onursatici]

I think the PR description is now stale, still refers to the fallback for >2G arrays

[onursatici]

currently we only support the varbin export to arrow, and don't go from a vortex varbin array, but I think that is fine for now

[github-actions]

Polar Signals Profiling Results

Latest Run

Status	Commit	Job	Attempt	Link
🟢 Done	2cc117e		1	Explore Profiling Data

---

Powered by Polar Signals Cloud