feat(cuda): add a FastLanes delta decode kernel

.github/workflows/codspeed.yml

@@ -256,7 +256,7 @@ jobs:
         include:
           - { shard: 1, name: "Bitpacked", benches: "bitpacked_cuda" }
           - { shard: 2, name: "Dynamic dispatch", benches: "dynamic_dispatch_cuda" }
-          - { shard: 3, name: "Standalone kernels", benches: "alp_cuda date_time_parts_cuda dict_cuda fsst_cuda runend_cuda" }
+          - { shard: 3, name: "Standalone kernels", benches: "alp_cuda date_time_parts_cuda delta_cuda dict_cuda fsst_cuda runend_cuda" }
     name: "Benchmark with Codspeed (CUDA Shard #${{ matrix.shard }} - ${{ matrix.name }})"
     timeout-minutes: 30
     runs-on: >-

vortex-btrblocks/src/builder.rs

@@ -178,8 +178,10 @@ impl BtrBlocksCompressorBuilder {
             string::StringDictScheme.id(),
             binary::BinaryDictScheme.id(),
         ];
-        // Delta has no GPU decode kernel and its prefix-sum decode is inherently sequential, so it
-        // is incompatible with pure-GPU decompression paths.
+        // Delta now has a CUDA decode kernel, so arrays that reach the GPU already encoded with
+        // it — the Delta children OnPair emits, for instance — decode there. It stays excluded
+        // from this preset until GPU delta decode is benchmarked against the schemes it would
+        // displace, since the preset picks encodings rather than merely decoding them.
         #[cfg(feature = "unstable_encodings")]
         excluded.push(integer::DeltaScheme::default().id());
         #[cfg(feature = "pco")]

vortex-cuda/Cargo.toml

@@ -65,6 +65,10 @@ fastlanes = { workspace = true }
 name = "for_cuda"
 harness = false
 
+[[bench]]
+name = "delta_cuda"
+harness = false
+
 [[bench]]
 name = "dict_cuda"
 harness = false

vortex-cuda/benches/delta_cuda.rs

@@ -0,0 +1,114 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright the Vortex contributors
+
+//! CUDA benchmarks for FastLanes delta decompression.
+//!
+//! Each element width is its own case: the lane count is `1024 / bit-width`, so the width sets
+//! both how a chunk is split (128 lanes of 8 rows for `u8` through 16 lanes of 64 rows for
+//! `u64`) and how much of a block is busy during the scan.
+
+#![expect(clippy::unwrap_used)]
+
+mod bench_config;
+mod timed_launch_strategy;
+
+use std::mem::size_of;
+use std::sync::Arc;
+use std::sync::atomic::Ordering;
+use std::time::Duration;
+
+use criterion::BenchmarkId;
+use criterion::Criterion;
+use criterion::Throughput;
+use cudarc::driver::DeviceRepr;
+use futures::executor::block_on;
+use vortex::array::IntoArray;
+use vortex::array::VortexSessionExecute;
+use vortex::array::array_session;
+use vortex::array::arrays::PrimitiveArray;
+use vortex::array::validity::Validity;
+use vortex::buffer::Buffer;
+use vortex::dtype::NativePType;
+use vortex::encodings::fastlanes::Delta;
+use vortex::encodings::fastlanes::DeltaArray;
+use vortex::error::VortexExpect;
+use vortex_cuda::CudaDispatchMode;
+use vortex_cuda::CudaSession;
+use vortex_cuda::executor::CudaArrayExt;
+use vortex_cuda_macros::cuda_available;
+use vortex_cuda_macros::cuda_not_available;
+
+use crate::bench_config::BENCH_SIZES;
+use crate::timed_launch_strategy::TimedLaunchStrategy;
+
+/// Builds a delta-encoded array of `len` values that stay inside `T`.
+fn make_delta_array<T>(len: usize) -> DeltaArray
+where
+    T: NativePType + From<u8>,
+{
+    // A small repeating step keeps every delta narrow, which is the shape delta is chosen for.
+    let data: Vec<T> = (0..len)
+        .map(|i| <T as From<u8>>::from(u8::try_from(i % 251).vortex_expect("modulo fits u8")))
+        .collect();
+    let primitive = PrimitiveArray::new(Buffer::from(data), Validity::NonNullable);
+
+    let mut ctx = array_session().create_execution_ctx();
+    Delta::try_from_primitive_array(&primitive, &mut ctx).vortex_expect("failed to delta encode")
+}
+
+fn benchmark_delta_typed<T>(c: &mut Criterion, type_name: &str)
+where
+    T: NativePType + DeviceRepr + From<u8>,
+{
+    let mut group = c.benchmark_group("cuda");
+
+    for &(len, len_str) in BENCH_SIZES {
+        group.throughput(Throughput::Bytes((len * size_of::<T>()) as u64));
+
+        let delta = make_delta_array::<T>(len);
+
+        group.bench_with_input(
+            BenchmarkId::new(format!("cuda/delta/{type_name}"), len_str),
+            &delta,
+            |b, delta| {
+                b.iter_custom(|iters| {
+                    let timed = TimedLaunchStrategy::default();
+                    let timer = timed.timer();
+
+                    let mut cuda_ctx =
+                        CudaSession::create_execution_ctx(&vortex_cuda::cuda_session())
+                            .vortex_expect("failed to create execution context")
+                            .with_dispatch_mode(CudaDispatchMode::StandaloneOnly)
+                            .with_launch_strategy(Arc::new(timed));
+
+                    for _ in 0..iters {
+                        block_on(delta.clone().into_array().execute_cuda(&mut cuda_ctx)).unwrap();
+                    }
+
+                    Duration::from_nanos(timer.load(Ordering::Relaxed))
+                });
+            },
+        );
+    }
+
+    group.finish();
+}
+
+fn benchmark_delta(c: &mut Criterion) {
+    benchmark_delta_typed::<u8>(c, "u8");
+    benchmark_delta_typed::<u16>(c, "u16");
+    benchmark_delta_typed::<u32>(c, "u32");
+    benchmark_delta_typed::<u64>(c, "u64");
+}
+
+criterion::criterion_group! {
+    name = benches;
+    config = bench_config::cuda_bench_config();
+    targets = benchmark_delta
+}
+
+#[cuda_available]
+criterion::criterion_main!(benches);
+
+#[cuda_not_available]
+fn main() {}

vortex-cuda/kernels/src/delta.cu

@@ -0,0 +1,78 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright the Vortex contributors
+
+#include "fastlanes_common.cuh"
+#include "types.cuh"
+
+// FastLanes delta decode.
+//
+// Delta is stored in the FastLanes transposed layout, where a 1024-element chunk is
+// FL_LANES<T> independent columns of (1024 / FL_LANES<T>) rows. Each column carries its own
+// running total seeded from that lane's base, so the chunk decodes as FL_LANES<T> independent
+// sequential scans rather than one scan over 1024 elements. One thread owns one lane.
+//
+// This mirrors `fastlanes::Delta::undelta` followed by `Transpose::untranspose`, which is what
+// the CPU decoder (`delta_decompress`) runs. Both steps must stay in step with that crate.
+
+/// Maps a position in the transposed layout to its position in natural order.
+///
+/// Mirrors `fastlanes::transpose`; `untranspose` is `output[transpose(i)] = input[i]`.
+__device__ inline uint32_t fl_transpose_index(uint32_t idx) {
+    const uint32_t lane = idx % 16;
+    const uint32_t order = (idx / 16) % 8;
+    const uint32_t row = idx / 128;
+    return (lane * 64) + (FL_ORDER[order] * 8) + row;
+}
+
+/// Decodes `num_chunks` full 1024-element delta chunks.
+///
+/// `deltas` and `output` hold `num_chunks * FL_CHUNK` elements; `bases` holds
+/// `num_chunks * FL_LANES<T>`. Only unsigned types are instantiated: the CPU decoder
+/// reinterprets signed input through its unsigned counterpart so that the wrapping add here
+/// inverts the wrapping subtract done at compress time.
+template <typename T>
+__device__ void delta_decode_kernel(const T *const __restrict deltas,
+                                    const T *const __restrict bases,
+                                    T *const __restrict output,
+                                    uint64_t num_chunks) {
+    constexpr uint32_t LANES = FL_LANES<T>;
+    constexpr uint32_t ROWS = FL_CHUNK / LANES;
+
+    // The undelta pass writes the chunk in transposed order; the untranspose pass then reads it
+    // back in a different order, so the whole chunk is staged in shared memory between them.
+    __shared__ T transposed[FL_CHUNK];
+
+    for (uint64_t chunk = blockIdx.x; chunk < num_chunks; chunk += gridDim.x) {
+        const T *const in = deltas + chunk * FL_CHUNK;
+        const T *const base = bases + chunk * LANES;
+
+        // Each lane accumulates down its own column, seeded by that lane's base.
+        for (uint32_t lane = threadIdx.x; lane < LANES; lane += blockDim.x) {
+            T running = base[lane];
+            for (uint32_t row = 0; row < ROWS; ++row) {
+                const uint32_t idx = INDEX(row, lane);
+                running = static_cast<T>(running + in[idx]);
+                transposed[idx] = running;
+            }
+        }
+        __syncthreads();
+
+        T *const out = output + chunk * FL_CHUNK;
+        for (uint32_t i = threadIdx.x; i < FL_CHUNK; i += blockDim.x) {
+            out[fl_transpose_index(i)] = transposed[i];
+        }
+        // Guard the staging buffer before the next chunk overwrites it.
+        __syncthreads();
+    }
+}
+
+#define GENERATE_DELTA_KERNEL(suffix, Type)                                                                  \
+    extern "C" __global__ void delta_##suffix(const Type *const __restrict deltas,                           \
+                                              const Type *const __restrict bases,                            \
+                                              Type *const __restrict output,                                 \
+                                              uint64_t num_chunks) {                                         \
+        delta_decode_kernel<Type>(deltas, bases, output, num_chunks);                                        \
+    }
+
+// Signed input is reinterpreted to its unsigned counterpart before the launch.
+FOR_EACH_UNSIGNED_INT(GENERATE_DELTA_KERNEL)