Add Arrow fixed-size-list ingress

cpp/src/interop/arrow_data_structures.cpp

@@ -21,6 +21,23 @@
 
 namespace cudf::interop {
 
+namespace {
+
+bool contains_fixed_size_list(ArrowSchema const& schema)
+{
+  ArrowSchemaView schema_view;
+  NANOARROW_THROW_NOT_OK(ArrowSchemaViewInit(&schema_view, &schema, nullptr));
+  if (schema_view.type == NANOARROW_TYPE_FIXED_SIZE_LIST) { return true; }
+
+  for (auto i = 0; i < schema.n_children; ++i) {
+    if (contains_fixed_size_list(*schema.children[i])) { return true; }
+  }
+
+  return schema.dictionary != nullptr && contains_fixed_size_list(*schema.dictionary);
+}
+
+}  // namespace
+
 /**
  * @brief A wrapper around ArrowDeviceArray data used for flexible lifetime management.
  *
@@ -66,6 +83,12 @@ struct arrow_array_container {
                         rmm::cuda_stream_view stream,
                         rmm::device_async_resource_ref mr)
   {
+    // TODO: Remove this guard when owning wrappers can export fixed-size-list with matching
+    // schema and buffers for every target device.
+    CUDF_EXPECTS(!contains_fixed_size_list(schema_),
+                 "Owning Arrow device wrappers do not support fixed-size-list input until "
+                 "fixed-size-list egress is implemented",
+                 cudf::data_type_error);
     switch (input_.device_type) {
       case ARROW_DEVICE_CUDA:
       case ARROW_DEVICE_CUDA_HOST:

cpp/src/interop/arrow_utilities.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -11,6 +11,9 @@
 
 #include <nanoarrow/nanoarrow.h>
 
+#include <limits>
+#include <stdexcept>
+
 namespace cudf {
 namespace detail {
 data_type arrow_to_cudf_type(ArrowSchemaView const* arrow_view)
@@ -33,7 +36,8 @@ data_type arrow_to_cudf_type(ArrowSchemaView const* arrow_view)
     case NANOARROW_TYPE_STRING_VIEW:
     case NANOARROW_TYPE_LARGE_STRING: return data_type(type_id::STRING);
     case NANOARROW_TYPE_LIST:
-    case NANOARROW_TYPE_LARGE_LIST: return data_type(type_id::LIST);
+    case NANOARROW_TYPE_LARGE_LIST:
+    case NANOARROW_TYPE_FIXED_SIZE_LIST: return data_type(type_id::LIST);
     case NANOARROW_TYPE_DICTIONARY: return data_type(type_id::DICTIONARY32);
     case NANOARROW_TYPE_STRUCT: return data_type(type_id::STRUCT);
     case NANOARROW_TYPE_TIMESTAMP: {
@@ -62,6 +66,61 @@ data_type arrow_to_cudf_type(ArrowSchemaView const* arrow_view)
   }
 }
 
+bool is_fixed_size_list(ArrowSchemaView const* arrow_view)
+{
+  return arrow_view->type == NANOARROW_TYPE_FIXED_SIZE_LIST;
+}
+
+size_type fixed_size_list_width(ArrowSchemaView const* arrow_view)
+{
+  CUDF_EXPECTS(
+    is_fixed_size_list(arrow_view), "Expected a fixed-size-list schema", cudf::data_type_error);
+  CUDF_EXPECTS(arrow_view->fixed_size >= 0,
+               "fixed-size-list width must be non-negative",
+               std::invalid_argument);
+  CUDF_EXPECTS(arrow_view->fixed_size <= std::numeric_limits<size_type>::max(),
+               "fixed-size-list width exceeds cuDF's maximum supported row count (cudf::size_type)",
+               std::overflow_error);
+  return static_cast<size_type>(arrow_view->fixed_size);
+}
+
+fixed_size_list_layout get_fixed_size_list_layout(ArrowSchemaView const* arrow_view,
+                                                  ArrowArray const* input)
+{
+  CUDF_EXPECTS(input->offset >= 0 && input->length >= 0,
+               "fixed-size-list offset and length must be non-negative",
+               std::invalid_argument);
+
+  constexpr auto max_size = static_cast<int64_t>(std::numeric_limits<size_type>::max());
+  CUDF_EXPECTS(input->length < max_size,
+               "fixed-size-list length exceeds cuDF's maximum supported row count "
+               "(cudf::size_type)",
+               std::overflow_error);
+  CUDF_EXPECTS(input->offset <= std::numeric_limits<int64_t>::max() - input->length,
+               "fixed-size-list row bounds overflow Arrow's int64 representation",
+               std::overflow_error);
+
+  auto const width    = fixed_size_list_width(arrow_view);
+  auto const num_rows = static_cast<size_type>(input->length);
+  auto const row_end  = input->offset + input->length;
+
+  // Width zero is valid for a foreign Arrow producer even though nanoarrow's schema builder
+  // rejects it. Its offsets and child bounds are all zero.
+  if (width == 0) { return {width, num_rows, input->offset, row_end, 0, 0, 0}; }
+
+  CUDF_EXPECTS(row_end <= std::numeric_limits<int64_t>::max() / width,
+               "fixed-size-list child bounds overflow Arrow's int64 representation",
+               std::overflow_error);
+  auto const child_length = input->length * width;
+  CUDF_EXPECTS(child_length <= max_size,
+               "Number of fixed-size-list child elements exceeds cuDF's maximum supported "
+               "row count (cudf::size_type)",
+               std::overflow_error);
+
+  return {
+    width, num_rows, input->offset, row_end, input->offset * width, child_length, row_end * width};
+}
+
 ArrowType id_to_arrow_type(cudf::type_id id)
 {
   switch (id) {

cpp/src/interop/arrow_utilities.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -28,6 +28,53 @@ static constexpr int fixed_width_data_buffer_idx = 1;
  */
 data_type arrow_to_cudf_type(ArrowSchemaView const* arrow_view);
 
+/**
+ * @brief Check whether the given schema view describes an Arrow fixed-size-list
+ *
+ * @param arrow_view SchemaView to check
+ * @return True if the schema describes a fixed-size-list
+ */
+bool is_fixed_size_list(ArrowSchemaView const* arrow_view);
+
+/**
+ * @brief Validated physical bounds for an Arrow fixed-size-list array
+ */
+struct fixed_size_list_layout {
+  size_type width;       ///< Child elements per row
+  size_type num_rows;    ///< Number of output rows
+  int64_t row_offset;    ///< First logical row in the Arrow array
+  int64_t row_end;       ///< One-past-last logical row
+  int64_t child_offset;  ///< First referenced child element
+  int64_t child_length;  ///< Number of referenced child elements
+  int64_t child_end;     ///< One-past-last referenced child element
+};
+
+/**
+ * @brief Return the number of child elements per row of a fixed-size-list schema
+ *
+ * @throw cudf::data_type_error if `arrow_view` is not a fixed-size-list
+ * @throw std::invalid_argument if the declared width is negative
+ * @throw std::overflow_error if the declared width exceeds `size_type`
+ *
+ * @param arrow_view SchemaView to pull the fixed size from
+ * @return Number of child elements per row
+ */
+size_type fixed_size_list_width(ArrowSchemaView const* arrow_view);
+
+/**
+ * @brief Validate and compute fixed-size-list row and child bounds
+ *
+ * @throw std::invalid_argument if row metadata is negative
+ * @throw std::overflow_error if Arrow bounds overflow `int64_t` or output lengths exceed
+ * `size_type`
+ *
+ * @param arrow_view Fixed-size-list schema view
+ * @param input Arrow array carrying row offset and length
+ * @return Validated source bounds and output sizes
+ */
+fixed_size_list_layout get_fixed_size_list_layout(ArrowSchemaView const* arrow_view,
+                                                  ArrowArray const* input);
+
 /**
  * @brief Map cudf column type id to ArrowType id
  *

cpp/src/interop/from_arrow_device.cu

@@ -4,6 +4,7 @@
  */
 
 #include "arrow_utilities.hpp"
+#include "from_arrow_host.hpp"
 
 #include <cudf/column/column_view.hpp>
 #include <cudf/copying.hpp>
@@ -321,15 +322,23 @@ dispatch_tuple_t dispatch_from_arrow_device::operator()<cudf::list_view>(
   CUDF_EXPECTS(schema->type != NANOARROW_TYPE_LARGE_LIST,
                "Large list types are not supported",
                cudf::data_type_error);
-  size_type const num_rows   = input->length;
-  size_type const offset     = input->offset;
+  auto const fixed_size = is_fixed_size_list(schema);
+  auto const layout =
+    fixed_size ? get_fixed_size_list_layout(schema, input) : fixed_size_list_layout{};
+
+  if (fixed_size) {
+    constexpr auto max_size = static_cast<int64_t>(std::numeric_limits<size_type>::max());
+    CUDF_EXPECTS(layout.row_end < max_size && layout.child_end <= max_size,
+                 "fixed-size-list device bounds exceed cuDF's maximum supported row count "
+                 "(cudf::size_type)",
+                 std::overflow_error);
+    CUDF_EXPECTS(input->children[0]->length >= layout.child_end,
+                 "fixed-size-list child is shorter than its parent layout requires",
+                 std::invalid_argument);
+  }
+  size_type const num_rows = fixed_size ? layout.num_rows : input->length;
+  size_type const offset   = fixed_size ? static_cast<size_type>(layout.row_offset) : input->offset;
   size_type const null_count = input->null_count;
-  auto offsets_view          = column_view{data_type(type_id::INT32),
-                                  (num_rows == 0) ? 0 : (offset + num_rows + 1),
-                                  input->buffers[fixed_width_data_buffer_idx],
-                                  nullptr,
-                                  0,
-                                  0};
 
   ArrowSchemaView child_schema_view;
   NANOARROW_THROW_NOT_OK(
@@ -341,8 +350,30 @@ dispatch_tuple_t dispatch_from_arrow_device::operator()<cudf::list_view>(
   // in the scenario where we were sliced and there are more elements in the child_view
   // than can be referenced by the sliced offsets, we need to slice the child_view
   // so that when `get_sliced_child` is called, we still produce the right result
-  auto max_child_offset =
-    num_rows == 0 ? 0 : cudf::detail::get_value<int32_t>(offsets_view, offset + num_rows, stream);
+  column_view offsets_view;
+  size_type max_child_offset = 0;
+  if (fixed_size) {
+    // fixed-size-list arrays carry no offsets buffer, so synthesize {0, w, 2w, ...}.
+    // these are absolute rather than normalized because the outer column_view applies a
+    // single offset to both the null mask and the children.
+    max_child_offset = (num_rows == 0) ? 0 : static_cast<size_type>(layout.child_end);
+    if (num_rows == 0) {
+      offsets_view = column_view{data_type{type_id::INT32}, 0, nullptr, nullptr, 0, 0};
+    } else {
+      owned.emplace_back(make_fixed_size_list_offsets(
+        static_cast<size_type>(layout.row_end) + 1, layout.width, stream, mr));
+      offsets_view = owned.back()->view();
+    }
+  } else {
+    offsets_view = column_view{data_type(type_id::INT32),
+                               (num_rows == 0) ? 0 : (offset + num_rows + 1),
+                               input->buffers[fixed_width_data_buffer_idx],
+                               nullptr,
+                               0,
+                               0};
+    max_child_offset =
+      num_rows == 0 ? 0 : cudf::detail::get_value<int32_t>(offsets_view, offset + num_rows, stream);
+  }
   child_view = cudf::slice(child_view, {0, max_child_offset}, stream).front();
 
   return std::make_tuple<column_view, owned_columns_t>(

cpp/src/interop/from_arrow_host.cu

@@ -30,6 +30,8 @@
 #include <rmm/device_buffer.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <thrust/sequence.h>
+
 #include <nanoarrow/nanoarrow.h>
 #include <nanoarrow/nanoarrow.hpp>
 #include <nanoarrow/nanoarrow_device.h>
@@ -284,24 +286,60 @@ std::unique_ptr<column> dispatch_copy_from_arrow_host::operator()<cudf::struct_v
     input->length, std::move(child_columns), null_count, std::move(*out_mask), stream, mr);
 }
 
+/**
+ * @brief Synthesize the offsets column and child bounds for a fixed-size-list array
+ *
+ * Mirrors the (offsets, child-offset, child-length) contract of `get_offsets_column`.
+ * Fixed-size-list arrays carry no offsets buffer, so `buffers[fixed_width_data_buffer_idx]`
+ * is never read here. The returned offsets are normalized to start at zero, matching
+ * `copy_offsets_column`; the absolute start of the child range is returned separately.
+ */
+std::tuple<std::unique_ptr<column>, int64_t, int64_t> get_fixed_size_list_offsets(
+  ArrowSchemaView const* schema,
+  ArrowArray const* input,
+  rmm::cuda_stream_view stream,
+  rmm::device_async_resource_ref mr)
+{
+  auto const layout = get_fixed_size_list_layout(schema, input);
+  CUDF_EXPECTS(input->children[0]->length >= layout.child_end,
+               "fixed-size-list child is shorter than its parent layout requires",
+               std::invalid_argument);
+
+  return std::tuple{make_fixed_size_list_offsets(layout.num_rows + 1, layout.width, stream, mr),
+                    layout.child_offset,
+                    layout.child_length};
+}
+
 template <>
 std::unique_ptr<column> dispatch_copy_from_arrow_host::operator()<cudf::list_view>(
   ArrowSchemaView const* schema, ArrowArray const* input, data_type type, bool skip_mask)
 {
+  CUDF_EXPECTS(input->length >= 0, "Number of rows must be non-negative.", std::invalid_argument);
   CUDF_EXPECTS(
-    input->length + 1 <= static_cast<std::int64_t>(std::numeric_limits<cudf::size_type>::max()),
+    input->length < static_cast<std::int64_t>(std::numeric_limits<cudf::size_type>::max()),
     "Number of rows exceeds cuDF's maximum supported row count (cudf::size_type).",
     std::overflow_error);
 
-  auto [offsets_column, offset, length] = get_offsets_column(schema, input, stream, mr);
+  auto const fixed_size                 = is_fixed_size_list(schema);
+  auto [offsets_column, offset, length] = fixed_size
+                                            ? get_fixed_size_list_offsets(schema, input, stream, mr)
+                                            : get_offsets_column(schema, input, stream, mr);
 
   ArrowSchemaView view;
   NANOARROW_THROW_NOT_OK(ArrowSchemaViewInit(&view, schema->schema->children[0], nullptr));
   auto child_type = arrow_to_cudf_type(&view);
 
   ArrowArray child_array(*input->children[0]);
+  if (fixed_size) {
+    CUDF_EXPECTS(child_array.offset >= 0,
+                 "fixed-size-list child offset must be non-negative",
+                 std::invalid_argument);
+    CUDF_EXPECTS(offset <= std::numeric_limits<int64_t>::max() - child_array.offset,
+                 "fixed-size-list child offset overflows Arrow's int64 representation",
+                 std::overflow_error);
+  }
   child_array.offset += offset;
-  child_array.length = std::min(length, child_array.length);
+  child_array.length = fixed_size ? length : std::min(length, child_array.length);
 
   auto child_column = get_column_copy(&view, &child_array, child_type, skip_mask, stream, mr);
 
@@ -394,9 +432,27 @@ std::tuple<std::unique_ptr<column>, int64_t, int64_t> copy_offsets_column(
 
 }  // namespace
 
+std::unique_ptr<column> make_fixed_size_list_offsets(size_type size,
+                                                     size_type width,
+                                                     rmm::cuda_stream_view stream,
+                                                     rmm::device_async_resource_ref mr)
+{
+  auto offsets =
+    make_numeric_column(data_type{type_id::INT32}, size, mask_state::UNALLOCATED, stream, mr);
+  auto d_offsets = offsets->mutable_view().begin<size_type>();
+  thrust::sequence(
+    rmm::exec_policy_nosync(stream, mr), d_offsets, d_offsets + size, size_type{0}, width);
+  return offsets;
+}
+
 /**
  * @brief Utility to copy the offsets from the given input (strings or list) to a
  * cudf column
+ *
+ * @note This requires `input` to carry an offsets buffer at
+ * `fixed_width_data_buffer_idx`, which it reads before inspecting `schema->type`.
+ * Fixed-size-list arrays have no offsets buffer (`n_buffers == 1`), so they must be
+ * routed to `get_fixed_size_list_offsets` instead.
  */
 std::tuple<std::unique_ptr<column>, int64_t, int64_t> get_offsets_column(
   ArrowSchemaView const* schema,

cpp/src/interop/from_arrow_host.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -48,5 +48,22 @@ std::tuple<std::unique_ptr<column>, int64_t, int64_t> get_offsets_column(
   rmm::cuda_stream_view stream,
   rmm::device_async_resource_ref mr);
 
+/**
+ * @brief Create the offsets column for a fixed-size-list array
+ *
+ * Arrow fixed-size-list arrays carry no offsets buffer; the offsets are implicit.
+ * This generates `size` offsets of the form `{0, width, 2*width, ...}`.
+ *
+ * @param size Number of offsets to generate (normally num_rows + 1)
+ * @param width Number of child elements per list row
+ * @param stream CUDA stream used for device memory operations
+ * @param mr Device memory resource to use for all device memory allocations
+ * @return INT32 offsets column
+ */
+std::unique_ptr<column> make_fixed_size_list_offsets(size_type size,
+                                                     size_type width,
+                                                     rmm::cuda_stream_view stream,
+                                                     rmm::device_async_resource_ref mr);
+
 }  // namespace detail
 }  // namespace cudf