Use owning Arrow types in C++ to expose data to Python

conda/environments/all_cuda-118_arch-aarch64.yaml

@@ -70,7 +70,7 @@ dependencies:
 - pyarrow>=14.0.0,<20.0.0a0
 - pydata-sphinx-theme>=0.15.4
 - pynvml>=12.0.0,<13.0.0a0
-- pytest-benchmark
+- pytest-benchmark<5.1.0
 - pytest-cases>=3.8.2
 - pytest-cov
 - pytest-rerunfailures

conda/environments/all_cuda-118_arch-x86_64.yaml

@@ -72,7 +72,7 @@ dependencies:
 - pyarrow>=14.0.0,<20.0.0a0
 - pydata-sphinx-theme>=0.15.4
 - pynvml>=12.0.0,<13.0.0a0
-- pytest-benchmark
+- pytest-benchmark<5.1.0
 - pytest-cases>=3.8.2
 - pytest-cov
 - pytest-rerunfailures

conda/environments/all_cuda-128_arch-aarch64.yaml

@@ -69,7 +69,7 @@ dependencies:
 - pydata-sphinx-theme>=0.15.4
 - pynvjitlink>=0.0.0a0
 - pynvml>=12.0.0,<13.0.0a0
-- pytest-benchmark
+- pytest-benchmark<5.1.0
 - pytest-cases>=3.8.2
 - pytest-cov
 - pytest-rerunfailures

conda/environments/all_cuda-128_arch-x86_64.yaml

@@ -70,7 +70,7 @@ dependencies:
 - pydata-sphinx-theme>=0.15.4
 - pynvjitlink>=0.0.0a0
 - pynvml>=12.0.0,<13.0.0a0
-- pytest-benchmark
+- pytest-benchmark<5.1.0
 - pytest-cases>=3.8.2
 - pytest-cov
 - pytest-rerunfailures

cpp/include/cudf/interop.hpp

@@ -314,7 +314,7 @@ class arrow_column {
    *
    * @param stream CUDA stream used for device memory operations and kernel launches
    * @param mr Device memory resource used for any allocations during conversion
-   * @return unique_column_view_t containing a view of the column data
+   * @return A view of the column data
    */
   [[nodiscard]] column_view view() const;
 
@@ -434,7 +434,7 @@ class arrow_table {
    *
    * @param stream CUDA stream used for device memory operations and kernel launches
    * @param mr Device memory resource used for any allocations during conversion
-   * @return unique_table_view_t containing a view of the table data
+   * @return A view of the table data
    */
   [[nodiscard]] table_view view() const;
 

cpp/src/column/column_view.cpp

@@ -37,7 +37,10 @@ template <typename ColumnView>
 void prefetch_col_data(ColumnView& col, void const* data_ptr, std::string_view key) noexcept
 {
   if (cudf::experimental::prefetch::detail::prefetch_config::instance().get(key)) {
-    if (cudf::is_fixed_width(col.type())) {
+    if (col.type().id() == cudf::type_id::EMPTY) {
+      // Skip prefetching for empty columns
+      return;
+    } else if (cudf::is_fixed_width(col.type())) {
       cudf::experimental::prefetch::detail::prefetch_noexcept(
         key, data_ptr, col.size() * size_of(col.type()), cudf::get_default_stream());
     } else if (col.type().id() == type_id::STRING) {

cpp/src/interop/from_arrow_device.cu

@@ -284,7 +284,7 @@ dispatch_tuple_t dispatch_from_arrow_device::operator()<cudf::list_view>(
   size_type const offset     = input->offset;
   size_type const null_count = input->null_count;
   auto offsets_view          = column_view{data_type(type_id::INT32),
-                                  offset + num_rows + 1,
+                                  (num_rows == 0) ? 0 : (offset + num_rows + 1),
                                   input->buffers[fixed_width_data_buffer_idx],
                                   nullptr,
                                   0,
@@ -300,8 +300,9 @@ 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 = cudf::detail::get_value<int32_t>(offsets_view, offset + num_rows, stream);
-  child_view            = cudf::slice(child_view, {0, max_child_offset}, stream).front();
+  auto 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>(
     {type,

cpp/src/interop/to_arrow_device.cu

@@ -91,6 +91,21 @@ int set_buffer(std::unique_ptr<T> device_buf, int64_t i, ArrowArray* out)
   return NANOARROW_OK;
 }
 
+int set_null_mask(column::contents& contents, ArrowArray* out)
+{
+  if (contents.null_mask) {
+    NANOARROW_RETURN_NOT_OK(set_buffer(std::move(contents.null_mask), validity_buffer_idx, out));
+  }
+  return NANOARROW_OK;
+}
+
+int set_contents(column::contents& contents, ArrowArray* out)
+{
+  NANOARROW_RETURN_NOT_OK(set_null_mask(contents, out));
+  NANOARROW_RETURN_NOT_OK(set_buffer(std::move(contents.data), fixed_width_data_buffer_idx, out));
+  return NANOARROW_OK;
+}
+
 struct dispatch_to_arrow_device {
   template <typename T,
             CUDF_ENABLE_IF(not is_rep_layout_compatible<T>() and not is_fixed_point<T>())>
@@ -117,23 +132,16 @@ struct dispatch_to_arrow_device {
     ArrowArrayMove(tmp.get(), out);
     return NANOARROW_OK;
   }
-
-  int set_null_mask(column::contents& contents, ArrowArray* out)
-  {
-    if (contents.null_mask) {
-      NANOARROW_RETURN_NOT_OK(set_buffer(std::move(contents.null_mask), validity_buffer_idx, out));
-    }
-    return NANOARROW_OK;
-  }
-
-  int set_contents(column::contents& contents, ArrowArray* out)
-  {
-    NANOARROW_RETURN_NOT_OK(set_null_mask(contents, out));
-    NANOARROW_RETURN_NOT_OK(set_buffer(std::move(contents.data), fixed_width_data_buffer_idx, out));
-    return NANOARROW_OK;
-  }
 };
 
+int handle_empty_type_column(ArrowArray* array, cudf::column& column)
+{
+  NANOARROW_RETURN_NOT_OK(initialize_array(array, NANOARROW_TYPE_NA, column.view()));
+  auto child_contents = column.release();
+  NANOARROW_RETURN_NOT_OK(set_contents(child_contents, array));
+  return NANOARROW_OK;
+}
+
 template <>
 int dispatch_to_arrow_device::operator()<bool>(cudf::column&& column,
                                                rmm::cuda_stream_view stream,
@@ -227,8 +235,12 @@ int dispatch_to_arrow_device::operator()<cudf::struct_view>(cudf::column&& colum
   for (size_t i = 0; i < size_t(tmp->n_children); ++i) {
     ArrowArray* child_ptr = tmp->children[i];
     auto& child           = contents.children[i];
-    NANOARROW_RETURN_NOT_OK(cudf::type_dispatcher(
-      child->type(), dispatch_to_arrow_device{}, std::move(*child), stream, mr, child_ptr));
+    if (child->type().id() == cudf::type_id::EMPTY) {
+      NANOARROW_RETURN_NOT_OK(handle_empty_type_column(child_ptr, *child));
+    } else {
+      NANOARROW_RETURN_NOT_OK(cudf::type_dispatcher(
+        child->type(), dispatch_to_arrow_device{}, std::move(*child), stream, mr, child_ptr));
+    }
   }
 
   ArrowArrayMove(tmp.get(), out);
@@ -253,8 +265,12 @@ int dispatch_to_arrow_device::operator()<cudf::list_view>(cudf::column&& column,
   NANOARROW_RETURN_NOT_OK(set_buffer(std::move(offsets_contents.data), 1, tmp.get()));
 
   auto& child = contents.children[cudf::lists_column_view::child_column_index];
-  NANOARROW_RETURN_NOT_OK(cudf::type_dispatcher(
-    child->type(), dispatch_to_arrow_device{}, std::move(*child), stream, mr, tmp->children[0]));
+  if (child->type().id() == cudf::type_id::EMPTY) {
+    NANOARROW_RETURN_NOT_OK(handle_empty_type_column(tmp->children[0], *child));
+  } else {
+    NANOARROW_RETURN_NOT_OK(cudf::type_dispatcher(
+      child->type(), dispatch_to_arrow_device{}, std::move(*child), stream, mr, tmp->children[0]));
+  }
 
   ArrowArrayMove(tmp.get(), out);
   return NANOARROW_OK;
@@ -535,8 +551,12 @@ unique_device_array_t to_arrow_device(cudf::table&& table,
   for (size_t i = 0; i < cols.size(); ++i) {
     auto child = tmp->children[i];
     auto col   = cols[i].get();
-    NANOARROW_THROW_NOT_OK(cudf::type_dispatcher(
-      col->type(), detail::dispatch_to_arrow_device{}, std::move(*col), stream, mr, child));
+    if (col->type().id() == cudf::type_id::EMPTY) {
+      NANOARROW_THROW_NOT_OK(handle_empty_type_column(child, *col));
+    } else {
+      NANOARROW_THROW_NOT_OK(cudf::type_dispatcher(
+        col->type(), detail::dispatch_to_arrow_device{}, std::move(*col), stream, mr, child));
+    }
   }
 
   return create_device_array(std::move(tmp), stream);
@@ -548,8 +568,12 @@ unique_device_array_t to_arrow_device(cudf::column&& col,
 {
   nanoarrow::UniqueArray tmp;
 
-  NANOARROW_THROW_NOT_OK(cudf::type_dispatcher(
-    col.type(), detail::dispatch_to_arrow_device{}, std::move(col), stream, mr, tmp.get()));
+  if (col.type().id() == cudf::type_id::EMPTY) {
+    NANOARROW_THROW_NOT_OK(handle_empty_type_column(tmp.get(), col));
+  } else {
+    NANOARROW_THROW_NOT_OK(cudf::type_dispatcher(
+      col.type(), detail::dispatch_to_arrow_device{}, std::move(col), stream, mr, tmp.get()));
+  }
 
   return create_device_array(std::move(tmp), stream);
 }

cpp/src/interop/to_arrow_schema.cpp

@@ -152,8 +152,12 @@ int dispatch_to_arrow_type::operator()<cudf::struct_view>(column_view input,
 
     child->flags = col.has_nulls() ? ARROW_FLAG_NULLABLE : 0;
 
-    NANOARROW_RETURN_NOT_OK(cudf::type_dispatcher(
-      col.type(), detail::dispatch_to_arrow_type{}, col, metadata.children_meta[i], child));
+    if (col.type().id() == cudf::type_id::EMPTY) {
+      NANOARROW_RETURN_NOT_OK(ArrowSchemaSetType(out->children[i], NANOARROW_TYPE_NA));
+    } else {
+      NANOARROW_RETURN_NOT_OK(cudf::type_dispatcher(
+        col.type(), detail::dispatch_to_arrow_type{}, col, metadata.children_meta[i], child));
+    }
   }
 
   return NANOARROW_OK;
@@ -174,6 +178,9 @@ int dispatch_to_arrow_type::operator()<cudf::list_view>(column_view input,
   out->flags = input.has_nulls() ? ARROW_FLAG_NULLABLE : 0;
   NANOARROW_RETURN_NOT_OK(ArrowSchemaSetName(out->children[0], child_meta.name.c_str()));
   out->children[0]->flags = child.has_nulls() ? ARROW_FLAG_NULLABLE : 0;
+  if (child.type().id() == cudf::type_id::EMPTY) {
+    return ArrowSchemaSetType(out->children[0], NANOARROW_TYPE_NA);
+  }
   return cudf::type_dispatcher(
     child.type(), detail::dispatch_to_arrow_type{}, child, child_meta, out->children[0]);
 }
@@ -219,8 +226,12 @@ unique_schema_t to_arrow_schema(cudf::table_view const& input,
     NANOARROW_THROW_NOT_OK(ArrowSchemaSetName(child, metadata[i].name.c_str()));
     child->flags = col.has_nulls() ? ARROW_FLAG_NULLABLE : 0;
 
-    NANOARROW_THROW_NOT_OK(
-      cudf::type_dispatcher(col.type(), detail::dispatch_to_arrow_type{}, col, metadata[i], child));
+    if (col.type().id() == cudf::type_id::EMPTY) {
+      NANOARROW_THROW_NOT_OK(ArrowSchemaSetType(child, NANOARROW_TYPE_NA));
+    } else {
+      NANOARROW_THROW_NOT_OK(cudf::type_dispatcher(
+        col.type(), detail::dispatch_to_arrow_type{}, col, metadata[i], child));
+    }
   }
 
   unique_schema_t out(new ArrowSchema, [](ArrowSchema* schema) {

dependencies.yaml

@@ -932,7 +932,14 @@ dependencies:
           - fastavro>=0.22.9
           - hypothesis
           - mmh3
-          - pytest-benchmark
+          # Version 5.1 is incompatible with pytest<8.2.
+          # https://github.com/ionelmc/pytest-benchmark/commit/8dfeeeca8a5c640a4dc4455904ace4f97f62e655
+          # Remove upper bound when we unbound pytest
+          # https://github.com/rapidsai/build-planning/issues/105
+          # or when a fixed version of pytest-benchmark 5.1 with the necessary bounds is released
+          # https://github.com/conda-forge/pytest-benchmark-feedstock/pull/27
+          # https://github.com/conda-forge/conda-forge-repodata-patches-feedstock/pull/990
+          - pytest-benchmark<5.1.0
           - pytest-cases>=3.8.2
           - scipy
           - zstandard

python/cudf/benchmarks/API/bench_dataframe.py

@@ -1,17 +1,24 @@
-# Copyright (c) 2022-2024, NVIDIA CORPORATION.
+# Copyright (c) 2022-2025, NVIDIA CORPORATION.
 
 """Benchmarks of DataFrame methods."""
 
 import string
 
 import numba.cuda
 import numpy
+import pyarrow as pa
 import pytest
 import pytest_cases
 from config import cudf, cupy
 from utils import benchmark_with_object
 
 
+@pytest.mark.parametrize("N", [100, 1_000_000, 100_000_000])
+def bench_from_arrow(benchmark, N):
+    rng = numpy.random.default_rng(seed=10)
+    benchmark(cudf.DataFrame, {None: pa.array(rng.random(N))})
+
+
 @pytest.mark.parametrize("N", [100, 1_000_000])
 def bench_construction(benchmark, N):
     benchmark(cudf.DataFrame, {None: cupy.random.rand(N)})

python/cudf/cudf/core/column/column.py

@@ -329,7 +329,7 @@ def null_count(self) -> int:
             else:
                 with acquire_spill_lock():
                     self._null_count = plc.null_mask.null_count(
-                        self.base_mask.get_ptr(mode="read"),  # type: ignore[union-attr]
+                        plc.gpumemoryview(self.base_mask),  # type: ignore[union-attr]
                         self.offset,
                         self.offset + self.size,
                     )
@@ -750,7 +750,7 @@ def to_arrow(self) -> pa.Array:
           4
         ]
         """
-        return plc.interop.to_arrow(self.to_pylibcudf(mode="read")).chunk(0)
+        return plc.interop.to_arrow(self.to_pylibcudf(mode="read"))
 
     @classmethod
     def from_arrow(cls, array: pa.Array) -> ColumnBase:

python/cudf/pyproject.toml

@@ -56,7 +56,7 @@ test = [
     "hypothesis",
     "mmh3",
     "msgpack",
-    "pytest-benchmark",
+    "pytest-benchmark<5.1.0",
     "pytest-cases>=3.8.2",
     "pytest-cov",
     "pytest-rerunfailures",

python/pylibcudf/pylibcudf/CMakeLists.txt

@@ -28,6 +28,7 @@ set(cython_sources
     groupby.pyx
     hashing.pyx
     interop.pyx
+    _interop_helpers.pyx
     jit.pyx
     join.pyx
     json.pyx
@@ -61,11 +62,13 @@ rapids_cython_create_modules(
   LINKED_LIBRARIES "${linked_libraries}" MODULE_PREFIX pylibcudf_ ASSOCIATED_TARGETS cudf
 )
 
-target_include_directories(pylibcudf_interop PUBLIC "$<BUILD_INTERFACE:${DLPACK_INCLUDE_DIR}>")
-
 include(${rapids-cmake-dir}/export/find_package_root.cmake)
 include(../../../cpp/cmake/thirdparty/get_nanoarrow.cmake)
-target_link_libraries(pylibcudf_interop PUBLIC nanoarrow)
+
+foreach(source interop _interop_helpers table column)
+  target_include_directories(pylibcudf_${source} PUBLIC "$<BUILD_INTERFACE:${DLPACK_INCLUDE_DIR}>")
+  target_link_libraries(pylibcudf_${source} PUBLIC nanoarrow)
+endforeach()
 
 add_subdirectory(libcudf)
 add_subdirectory(strings)

python/pylibcudf/pylibcudf/_interop_helpers.pxd

@@ -0,0 +1,9 @@
+# Copyright (c) 2025, NVIDIA CORPORATION.
+
+from pylibcudf.libcudf.interop cimport column_metadata
+
+cdef void _release_schema(object schema_capsule) noexcept
+
+cdef void _release_array(object array_capsule) noexcept
+
+cdef column_metadata _metadata_to_libcudf(metadata)