bench: run the full compress suite under --gpu-decompress

benchmarks/compress-bench/README.md

@@ -117,7 +117,11 @@ rendered before the failure summary, so a dataset the GPU cannot decode still le
 the matrix with numbers — the process exits non-zero either way.
 
 The dataset list in `src/main.rs` therefore holds only datasets a `--gpu-verify` run has confirmed.
-Several others are waiting on `vortex-cuda` kernel gaps (`u16` in `date_time_parts`, a
-`vortex.masked` kernel, and a CPU fallback reached with device-resident buffers); they are listed
-with their reasons next to `gpu_datasets`. Add one there once its gap is closed and verification
-passes.
+It now covers the whole compress suite: the kernel gaps that kept `taxi`, `Arade`, `CMSprovider`,
+`Euro2016`, `HashTags` and the `StructListOfInts` wide tables off it — `u16` components in
+`date_time_parts`, per-element `RunEnd` validity, and missing `vortex.masked` and `vortex.list`
+kernels — have since been closed. Add a new dataset there once verification passes.
+
+`airquality` and `rplace` download from pcodec's public bucket, which the CPU suite skips to avoid
+creating egress charges for pcodec. The GPU suite runs every entry on its explicit list, so both
+are fetched on each GPU run.

benchmarks/compress-bench/src/main.rs

@@ -234,24 +234,23 @@ async fn run_compress(
 
     // Datasets run in GPU mode. Add one only after a `--gpu-verify` run has confirmed its CUDA
     // decode end to end; a dataset here that cannot decode fails the benchmark job. Between them
-    // these cover FSST strings, bit-packed numerics and columns with nulls.
-    //
-    // Not yet listed, each blocked on a `vortex-cuda` gap rather than on the benchmark:
-    //
-    // - `taxi` and `Arade`: `Unsupported ptype u16`. The CUDA `date_time_parts` kernel dispatches
-    //   with `match_each_signed_integer_ptype!` where the CPU canonicaliser uses
-    //   `match_each_integer_ptype!`. Widening the fused kernel takes 4³ = 64 PTX instantiations
-    //   to 8³ = 512, so it is not a free change.
-    // - `Euro2016` and `HashTags`: `No CUDA kernel for encoding vortex.masked`.
-    // - `CMSprovider`: `expected host buffer` — a CPU fallback is reached with device-resident
-    //   buffers, which `CudaArrayExt::execute_cuda` refuses.
-    // - `StructListOfInts`: its list layouts have no verified CUDA decode path.
-    let gpu_datasets: [&dyn Dataset; 4] = [
+    // these cover FSST strings, bit-packed numerics, timestamps, columns with nulls, and lists.
+    let gpu_datasets: Vec<&dyn Dataset> = [
         &TPCHLCommentCanonical as &dyn Dataset,
         &TPCHLCommentChunked,
+        &TaxiData,
+        PBI_DATASETS.get(Arade),
         PBI_DATASETS.get(Bimbo),
+        PBI_DATASETS.get(CMSprovider),
+        PBI_DATASETS.get(Euro2016),
         PBI_DATASETS.get(Food),
-    ];
+        PBI_DATASETS.get(HashTags),
+        &DownloadableDataset::AirQuality,
+        &DownloadableDataset::RPlace,
+    ]
+    .into_iter()
+    .chain(structlistofints.iter().map(|d| d as &dyn Dataset))
+    .collect();
 
     let all_datasets: Vec<&dyn Dataset> = [
         &TaxiData as &dyn Dataset,
@@ -276,14 +275,17 @@ async fn run_compress(
     .collect();
 
     let datasets: Vec<&dyn Dataset> = if mode.is_gpu() {
-        gpu_datasets.to_vec()
+        gpu_datasets
     } else {
         all_datasets
     }
     .into_iter()
     .filter(|d| {
         if let Some(filter) = datasets_filter.as_ref() {
             filter.is_match(d.name())
+        } else if mode.is_gpu() {
+            // The GPU suite is an explicit list, so every entry on it runs.
+            true
         } else {
             // These download data from pcodec's public bucket, presumably creating egress charges
             // for pcodec. As such, we do not run in CI.

vortex-btrblocks/src/builder.rs

@@ -175,6 +175,7 @@ impl BtrBlocksCompressorBuilder {
             float::ALPRDScheme.id(),
             float::FloatRLEScheme.id(),
             float::NullDominatedSparseScheme.id(),
+            string::NullDominatedSparseScheme.id(),
             string::StringDictScheme.id(),
             binary::BinaryDictScheme.id(),
         ];
@@ -268,6 +269,22 @@ mod tests {
         );
     }
 
+    /// `vortex.sparse` has no CUDA decode kernel, so no sparse scheme may survive this preset.
+    #[test]
+    fn cuda_compatible_excludes_every_sparse_scheme() {
+        let builder = BtrBlocksCompressorBuilder::default().only_cuda_compatible();
+        for excluded in [
+            integer::SparseScheme.id(),
+            float::NullDominatedSparseScheme.id(),
+            string::NullDominatedSparseScheme.id(),
+        ] {
+            assert!(
+                !builder.schemes.iter().any(|s| s.id() == excluded),
+                "{excluded} should be excluded"
+            );
+        }
+    }
+
     #[test]
     fn cuda_compatible_uses_fsst_for_strings() {
         let builder = BtrBlocksCompressorBuilder::default().only_cuda_compatible();

vortex-cuda/kernels/src/date_time_parts.cu

@@ -43,22 +43,34 @@ __device__ void date_time_parts(const DaysT *__restrict days,
     }
 
 #define EXPAND_DAYS(X)                                                                                       \
+    X(u8, uint8_t)                                                                                           \
+    X(u16, uint16_t)                                                                                         \
+    X(u32, uint32_t)                                                                                         \
+    X(u64, uint64_t)                                                                                         \
     X(i8, int8_t)                                                                                            \
     X(i16, int16_t)                                                                                          \
     X(i32, int32_t)                                                                                          \
     X(i64, int64_t)
 
 #define EXPAND_SUBSECONDS(d, DT, s, ST)                                                                      \
+    GENERATE_DATE_TIME_PARTS_KERNEL(d, DT, s, ST, u8, uint8_t)                                               \
+    GENERATE_DATE_TIME_PARTS_KERNEL(d, DT, s, ST, u16, uint16_t)                                             \
+    GENERATE_DATE_TIME_PARTS_KERNEL(d, DT, s, ST, u32, uint32_t)                                             \
+    GENERATE_DATE_TIME_PARTS_KERNEL(d, DT, s, ST, u64, uint64_t)                                             \
     GENERATE_DATE_TIME_PARTS_KERNEL(d, DT, s, ST, i8, int8_t)                                                \
     GENERATE_DATE_TIME_PARTS_KERNEL(d, DT, s, ST, i16, int16_t)                                              \
     GENERATE_DATE_TIME_PARTS_KERNEL(d, DT, s, ST, i32, int32_t)                                              \
     GENERATE_DATE_TIME_PARTS_KERNEL(d, DT, s, ST, i64, int64_t)
 
 #define EXPAND_SECONDS(d, DT)                                                                                \
+    EXPAND_SUBSECONDS(d, DT, u8, uint8_t)                                                                    \
+    EXPAND_SUBSECONDS(d, DT, u16, uint16_t)                                                                  \
+    EXPAND_SUBSECONDS(d, DT, u32, uint32_t)                                                                  \
+    EXPAND_SUBSECONDS(d, DT, u64, uint64_t)                                                                  \
     EXPAND_SUBSECONDS(d, DT, i8, int8_t)                                                                     \
     EXPAND_SUBSECONDS(d, DT, i16, int16_t)                                                                   \
     EXPAND_SUBSECONDS(d, DT, i32, int32_t)                                                                   \
     EXPAND_SUBSECONDS(d, DT, i64, int64_t)
 
-// Generate all 64 kernels (4³)
+// Generate all 512 kernels (8³: every signed and unsigned integer width per component)
 EXPAND_DAYS(EXPAND_SECONDS)

vortex-cuda/kernels/src/list.cu

@@ -0,0 +1,36 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright the Vortex contributors
+
+#include "config.cuh"
+#include "types.cuh"
+
+// Converts Arrow-style `List` offsets into `ListView` offset/size pairs.
+//
+// `List` stores `list_len + 1` monotonically increasing offsets; a `ListView` stores one offset
+// and one size per list. Both outputs are written by the same thread so the two views of a list
+// are always produced together.
+template <typename OffsetT>
+__device__ void list_views(const OffsetT *const __restrict offsets,
+                           OffsetT *const __restrict out_offsets,
+                           OffsetT *const __restrict out_sizes,
+                           uint64_t list_len) {
+    const uint32_t elements_per_block = blockDim.x * ELEMENTS_PER_THREAD;
+    const uint64_t block_start = static_cast<uint64_t>(blockIdx.x) * elements_per_block;
+    const uint64_t block_end = min(block_start + elements_per_block, list_len);
+
+    for (uint64_t idx = block_start + threadIdx.x; idx < block_end; idx += blockDim.x) {
+        const OffsetT start = offsets[idx];
+        out_offsets[idx] = start;
+        out_sizes[idx] = static_cast<OffsetT>(offsets[idx + 1] - start);
+    }
+}
+
+#define GENERATE_LIST_VIEWS_KERNEL(offset_suffix, OffsetT)                                                   \
+    extern "C" __global__ void list_views_##offset_suffix(const OffsetT *const __restrict offsets,           \
+                                                          OffsetT *const __restrict out_offsets,             \
+                                                          OffsetT *const __restrict out_sizes,               \
+                                                          uint64_t list_len) {                               \
+        list_views<OffsetT>(offsets, out_offsets, out_sizes, list_len);                                      \
+    }
+
+FOR_EACH_INTEGER(GENERATE_LIST_VIEWS_KERNEL)

vortex-cuda/kernels/src/runend.cu

@@ -127,6 +127,45 @@ __device__ void runend_decode_kernel(const EndsT *const __restrict ends,
     }
 }
 
+// Expands run-end encoded validity bits into a packed output bitmap.
+//
+// Mirrors `runend_decode_kernel`, but each thread owns one complete output byte so that
+// threads never race on bits within the same byte. Runs are located with a global binary
+// search per element rather than the shared-memory cache: validity expansion runs once per
+// array and is not the decode hot path.
+template <typename EndsT>
+__device__ void runend_bool_kernel(const EndsT *const __restrict ends,
+                                   uint64_t num_runs,
+                                   const uint8_t *const __restrict values,
+                                   uint64_t values_bit_offset,
+                                   uint64_t offset,
+                                   uint64_t output_len,
+                                   uint8_t *const __restrict output) {
+    const uint64_t output_bytes = (output_len + 7) / 8;
+    const uint32_t elements_per_block = blockDim.x * ELEMENTS_PER_THREAD;
+    const uint64_t block_start = static_cast<uint64_t>(blockIdx.x) * elements_per_block;
+    const uint64_t block_end = min(block_start + elements_per_block, output_bytes);
+
+    for (uint64_t byte_idx = block_start + threadIdx.x; byte_idx < block_end; byte_idx += blockDim.x) {
+        const uint64_t row_start = byte_idx * 8;
+        uint8_t packed = 0;
+#pragma unroll
+        for (uint32_t bit = 0; bit < 8; ++bit) {
+            const uint64_t row = row_start + bit;
+            if (row < output_len) {
+                uint64_t run_idx = upper_bound(ends, num_runs, row + offset);
+                if (run_idx >= num_runs) {
+                    run_idx = num_runs - 1;
+                }
+                const uint64_t value_idx = values_bit_offset + run_idx;
+                const uint8_t value = (values[value_idx / 8] >> (value_idx % 8)) & 1;
+                packed |= static_cast<uint8_t>(value << bit);
+            }
+        }
+        output[byte_idx] = packed;
+    }
+}
+
 #define GENERATE_RUNEND_KERNEL(value_suffix, ValueType, ends_suffix, EndsType)                               \
     extern "C" __global__ void runend_##value_suffix##_##ends_suffix(                                        \
         const EndsType *const __restrict ends,                                                               \
@@ -155,3 +194,21 @@ GENERATE_RUNEND_KERNELS_FOR_VALUE(i64, int64_t)
 GENERATE_RUNEND_KERNELS_FOR_VALUE(f16, __half)
 GENERATE_RUNEND_KERNELS_FOR_VALUE(f32, float)
 GENERATE_RUNEND_KERNELS_FOR_VALUE(f64, double)
+
+#define GENERATE_RUNEND_BOOL_KERNEL(ends_suffix, EndsType)                                                   \
+    extern "C" __global__ void runend_bool_##ends_suffix(const EndsType *const __restrict ends,              \
+                                                         uint64_t num_runs,                                  \
+                                                         const uint8_t *const __restrict values,             \
+                                                         uint64_t values_bit_offset,                         \
+                                                         uint64_t offset,                                    \
+                                                         uint64_t output_len,                                \
+                                                         uint8_t *const __restrict output) {                 \
+        runend_bool_kernel<EndsType>(ends, num_runs, values, values_bit_offset, offset, output_len, output); \
+    }
+
+// Validity bitmaps use a different physical layout and launch unit, but dispatch over the
+// same run-end index types.
+GENERATE_RUNEND_BOOL_KERNEL(u8, uint8_t)
+GENERATE_RUNEND_BOOL_KERNEL(u16, uint16_t)
+GENERATE_RUNEND_BOOL_KERNEL(u32, uint32_t)
+GENERATE_RUNEND_BOOL_KERNEL(u64, uint64_t)

vortex-cuda/src/canonical.rs

@@ -5,18 +5,21 @@ use std::sync::Arc;
 
 use async_trait::async_trait;
 use futures::future::try_join_all;
+use vortex::array::ArrayRef;
 use vortex::array::Canonical;
 use vortex::array::IntoArray;
 use vortex::array::VortexSessionExecute;
 use vortex::array::arrays::BoolArray;
 use vortex::array::arrays::DecimalArray;
 use vortex::array::arrays::ExtensionArray;
+use vortex::array::arrays::ListViewArray;
 use vortex::array::arrays::PrimitiveArray;
 use vortex::array::arrays::StructArray;
 use vortex::array::arrays::VarBinViewArray;
 use vortex::array::arrays::bool::BoolDataParts;
 use vortex::array::arrays::decimal::DecimalDataParts;
 use vortex::array::arrays::extension::ExtensionArrayExt;
+use vortex::array::arrays::listview::ListViewDataParts;
 use vortex::array::arrays::primitive::PrimitiveDataParts;
 use vortex::array::arrays::struct_::StructDataParts;
 use vortex::array::arrays::varbinview::BinaryView;
@@ -29,6 +32,19 @@ use vortex::buffer::Buffer;
 use vortex::buffer::ByteBuffer;
 use vortex::error::VortexResult;
 
+/// Copy a canonical child array to the host.
+async fn child_into_host(child: ArrayRef) -> VortexResult<ArrayRef> {
+    #[expect(
+        clippy::disallowed_methods,
+        reason = "CanonicalCudaExt threads no session through"
+    )]
+    Ok(child
+        .execute::<Canonical>(&mut legacy_session().create_execution_ctx())?
+        .into_host()
+        .await?
+        .into_array())
+}
+
 /// Move all canonical data from to_host from device.
 #[async_trait]
 pub trait CanonicalCudaExt {
@@ -102,18 +118,15 @@ impl CanonicalCudaExt for Canonical {
             n @ Canonical::Null(_) => Ok(n),
             Canonical::Bool(bool) => {
                 let len = bool.len();
-                let validity = bool.validity()?;
+                let validity = validity_into_host(bool.validity()?).await?;
                 let BoolDataParts { bits, meta } = bool.into_data().into_parts(len);
 
                 let bits = BitBuffer::new_with_offset(
                     bits.try_into_host()?.await?,
                     meta.len(),
                     meta.offset(),
                 );
-                Ok(Canonical::Bool(BoolArray::new(
-                    bits,
-                    validity_into_host(validity).await?,
-                )))
+                Ok(Canonical::Bool(BoolArray::new(bits, validity)))
             }
             Canonical::Primitive(prim) => {
                 let PrimitiveDataParts {
@@ -153,6 +166,7 @@ impl CanonicalCudaExt for Canonical {
                     validity,
                     dtype,
                 } = varbinview.into_data_parts();
+                let validity = validity_into_host(validity).await?;
 
                 // Copy all device views to host
                 let host_views = views.try_into_host()?.await?;
@@ -167,11 +181,27 @@ impl CanonicalCudaExt for Canonical {
                 let host_buffers = try_join_all(host_buffers).await?;
                 let host_buffers: Arc<[ByteBuffer]> = Arc::from(host_buffers);
 
-                let validity = validity_into_host(validity).await?;
                 Ok(Canonical::VarBinView(unsafe {
                     VarBinViewArray::new_unchecked(host_views, host_buffers, dtype, validity)
                 }))
             }
+            Canonical::List(list) => {
+                let ListViewDataParts {
+                    elements,
+                    offsets,
+                    sizes,
+                    validity,
+                    ..
+                } = list.into_data_parts();
+                let validity = validity_into_host(validity).await?;
+
+                Ok(Canonical::List(ListViewArray::try_new(
+                    child_into_host(elements).await?,
+                    child_into_host(offsets).await?,
+                    child_into_host(sizes).await?,
+                    validity,
+                )?))
+            }
             Canonical::Extension(ext) => {
                 // Copy the storage array to host and rewrap in ExtensionArray.
                 #[expect(