clean: simplify Arrow device export internals

Cleanups from a review of the Arrow device export path, no behavior
change:

- Inline the single-caller list_layout_parts 5-tuple helper and the
  export_list_layout_with_child indirection into export_list and
  export_list_layout. ListChildExport itself stays: the take-based host
  list-view rebuild wraps children in transient Dict encodings that the
  pre-computed schema does not include, so the rebuilt-child arm must
  keep canonicalizing; its doc now spells that out.
- Remove the dead end < offset wrap check in arrow_binary.cu: both
  addends are u32 widened to u64, so the sum cannot wrap.
- Replace the panicking as_ptype on the dyn-dispatch cast target with
  PType::try_from, matching the source-side check three lines up.
- Correct the export_binary rationale: cuDF's Arrow Device import
  rejects both Binary and BinaryView today (arrow_to_cudf_type maps
  neither), matching the e2e harness comment.
- Fix the misplaced doc comment on export_list_view, and document the
  validity buffer padding contract on export_arrow_validity_buffer
  (every path is backed by a 4-byte-padded, tail-zeroed allocation via
  the device copy layer).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Alexander Droste <alexander.droste@protonmail.com>

Author: 0ax1

vortex-cuda/kernels/src/arrow_binary.cu

@@ -63,8 +63,9 @@ __device__ void init_scan_device(const BinaryView *const __restrict views,
             const BinaryViewRef *const view_ref = reinterpret_cast<const BinaryViewRef *>(&view);
             const uint64_t buffer_index = static_cast<uint64_t>(view_ref->buffer_index);
             const uint64_t offset = static_cast<uint64_t>(view_ref->offset);
+            // Both addends are u32 widened to u64, so the end position cannot wrap.
             const uint64_t end = offset + static_cast<uint64_t>(size);
-            if (buffer_index >= data_buffer_count || end < offset || end > data_buffer_lens[buffer_index]) {
+            if (buffer_index >= data_buffer_count || end > data_buffer_lens[buffer_index]) {
                 scan[idx] = 0;
                 atomicMax(status, 1u);
                 continue;

vortex-cuda/src/arrow/canonical.rs

@@ -448,9 +448,11 @@ where
 
 /// Export Vortex binary views as standard Arrow `Binary`.
 ///
-/// cuDF imports Arrow `Binary` through the Arrow Device path, but does not currently accept
-/// Arrow `BinaryView`. This path keeps conversion on the CUDA stream by building `i32` offsets
-/// from view sizes and gathering inline/out-of-line view bytes into one contiguous values buffer.
+/// cuDF's Arrow Device import currently rejects both Arrow `Binary` and Arrow `BinaryView`
+/// (unlike `Utf8View`, which it accepts as strings), so standard `Binary` is exported as the
+/// layout other Arrow Device consumers accept most widely. This path keeps conversion on the
+/// CUDA stream by building `i32` offsets from view sizes and gathering inline/out-of-line view
+/// bytes into one contiguous values buffer.
 async fn export_binary(
     varbinview: VarBinViewArray,
     ctx: &mut CudaExecutionCtx,
@@ -679,6 +681,11 @@ fn gather_binary_values(
 /// Export Vortex validity as an Arrow validity byte buffer on the CUDA device.
 ///
 /// Returns `None` for the buffer when Arrow can omit validity because all rows are valid.
+///
+/// Every returned buffer is backed by an allocation padded to a 4-byte multiple with zeroed
+/// padding so cuDF's word-sized mask reads stay in bounds: the fast path through the device
+/// copy's tail zeroing, the other paths through their own padded allocations. Bits at positions
+/// `>= len + arrow_offset` within the final data byte are unspecified, as Arrow permits.
 pub(super) async fn export_arrow_validity_buffer(
     validity: Validity,
     len: usize,
@@ -788,7 +795,7 @@ pub(super) fn repack_arrow_validity_buffer(
     Ok(BufferHandle::new_device(Arc::new(output_device)).slice(0..output_bytes))
 }
 
-/// Export a standard Vortex list as Arrow `List`: validity, offsets, and one child array.
+/// Export a Vortex list-view as Arrow `List`.
 async fn export_list_view(
     listview: ListViewArray,
     ctx: &mut CudaExecutionCtx,
@@ -815,29 +822,12 @@ async fn export_list_view(
     .await
 }
 
+/// Export a standard Vortex list as Arrow `List`: validity, offsets, and one child array.
 async fn export_list(
     array: ListArray,
     child_export: ListChildExport,
     ctx: &mut CudaExecutionCtx,
 ) -> VortexResult<(ArrowArray, SyncEvent)> {
-    let (elements, len, validity_buffer, null_count, offsets_buffer) =
-        list_layout_parts(array, ctx).await?;
-    export_list_layout(
-        elements,
-        len,
-        validity_buffer,
-        null_count,
-        offsets_buffer,
-        child_export,
-        ctx,
-    )
-    .await
-}
-
-async fn list_layout_parts(
-    array: ListArray,
-    ctx: &mut CudaExecutionCtx,
-) -> VortexResult<(ArrayRef, usize, Option<BufferHandle>, i64, BufferHandle)> {
     let len = array.len();
     let ListDataParts {
         elements,
@@ -848,7 +838,16 @@ async fn list_layout_parts(
 
     let (validity_buffer, null_count) = export_arrow_validity_buffer(validity, len, 0, ctx).await?;
     let offsets_buffer = export_arrow_list_offsets(offsets, ctx).await?;
-    Ok((elements, len, validity_buffer, null_count, offsets_buffer))
+    export_list_layout(
+        elements,
+        len,
+        validity_buffer,
+        null_count,
+        offsets_buffer,
+        child_export,
+        ctx,
+    )
+    .await
 }
 
 #[derive(Clone, Copy)]
@@ -858,6 +857,11 @@ pub(super) enum ListChildExport {
     PreserveConcreteLayout,
     /// Canonicalize temporary encodings introduced by the host ListView
     /// rebuild, while still preserving rebuilt dictionary children.
+    ///
+    /// This is not equivalent to `export_array`: the take-based rebuild wraps
+    /// children in transient encodings (for example `take` returns `Dict`
+    /// arrays nested inside struct fields) that the pre-computed export schema
+    /// does not include, so they must be canonicalized away before export.
     RebuiltListViewChild,
 }
 
@@ -891,24 +895,7 @@ pub(super) async fn export_list_layout(
     ctx: &mut CudaExecutionCtx,
 ) -> VortexResult<(ArrowArray, SyncEvent)> {
     let elements_child = child_export.export(elements, ctx).await?;
-    export_list_layout_with_child(
-        elements_child,
-        len,
-        validity_buffer,
-        null_count,
-        offsets_buffer,
-        ctx,
-    )
-}
 
-fn export_list_layout_with_child(
-    elements_child: ArrowArray,
-    len: usize,
-    validity_buffer: Option<BufferHandle>,
-    null_count: i64,
-    offsets_buffer: BufferHandle,
-    ctx: &mut CudaExecutionCtx,
-) -> VortexResult<(ArrowArray, SyncEvent)> {
     let mut private_data = PrivateData::new(
         vec![validity_buffer, Some(offsets_buffer)],
         vec![elements_child],

vortex-cuda/src/dynamic_dispatch/plan_builder.rs

@@ -139,7 +139,9 @@ fn is_dyn_dispatch_cast_compatible(array: &ArrayRef) -> bool {
     let Ok(source_ptype) = PType::try_from(cast.child_at(0).dtype()) else {
         return false;
     };
-    let target_ptype = cast.scalar_fn().as_::<Cast>().as_ptype();
+    let Ok(target_ptype) = PType::try_from(cast.scalar_fn().as_::<Cast>()) else {
+        return false;
+    };
 
     // Implemented as unsigned dictionary-code casts to cuDF's signed index types.
     // LOAD/BITUNPACK materialize directly into the target-width output type.