Merge branch 'branch-24.12' into numpy_random

cpp/src/interop/to_arrow_host.cu

@@ -44,6 +44,7 @@
 #include <nanoarrow/nanoarrow.h>
 #include <nanoarrow/nanoarrow.hpp>
 #include <nanoarrow/nanoarrow_device.h>
+#include <sys/mman.h>
 
 #include <iostream>
 
@@ -52,6 +53,30 @@ namespace detail {
 
 namespace {
 
+/*
+  Enable Transparent Huge Pages (THP) for large (>4MB) allocations.
+  `buf` is returned untouched.
+  Enabling THP can improve performance of device-host memory transfers
+  significantly, see <https://github.com/rapidsai/cudf/pull/13914>.
+*/
+void enable_hugepage(ArrowBuffer* buffer)
+{
+  if (buffer->size_bytes < (1u << 22u)) {  // Smaller than 4 MB
+    return;
+  }
+
+#ifdef MADV_HUGEPAGE
+  auto const pagesize = sysconf(_SC_PAGESIZE);
+  void* addr          = const_cast<uint8_t*>(buffer->data);
+  auto length{static_cast<std::size_t>(buffer->size_bytes)};
+  if (std::align(pagesize, pagesize, addr, length)) {
+    // Intentionally not checking for errors that may be returned by older kernel versions;
+    // optimistically tries enabling huge pages.
+    madvise(addr, length, MADV_HUGEPAGE);
+  }
+#endif
+}
+
 struct dispatch_to_arrow_host {
   cudf::column_view column;
   rmm::cuda_stream_view stream;
@@ -62,6 +87,7 @@ struct dispatch_to_arrow_host {
     if (!column.has_nulls()) { return NANOARROW_OK; }
 
     NANOARROW_RETURN_NOT_OK(ArrowBitmapResize(bitmap, static_cast<int64_t>(column.size()), 0));
+    enable_hugepage(&bitmap->buffer);
     CUDF_CUDA_TRY(cudaMemcpyAsync(bitmap->buffer.data,
                                   (column.offset() > 0)
                                     ? cudf::detail::copy_bitmask(column, stream, mr).data()
@@ -76,6 +102,7 @@ struct dispatch_to_arrow_host {
   int populate_data_buffer(device_span<T const> input, ArrowBuffer* buffer) const
   {
     NANOARROW_RETURN_NOT_OK(ArrowBufferResize(buffer, input.size_bytes(), 1));
+    enable_hugepage(buffer);
     CUDF_CUDA_TRY(cudaMemcpyAsync(
       buffer->data, input.data(), input.size_bytes(), cudaMemcpyDefault, stream.value()));
     return NANOARROW_OK;

python/pylibcudf/pylibcudf/binaryop.pyx

@@ -52,33 +52,27 @@ cpdef Column binary_operation(
 
     if LeftBinaryOperand is Column and RightBinaryOperand is Column:
         with nogil:
-            result = move(
-                cpp_binaryop.binary_operation(
-                    lhs.view(),
-                    rhs.view(),
-                    op,
-                    output_type.c_obj
-                )
+            result = cpp_binaryop.binary_operation(
+                lhs.view(),
+                rhs.view(),
+                op,
+                output_type.c_obj
             )
     elif LeftBinaryOperand is Column and RightBinaryOperand is Scalar:
         with nogil:
-            result = move(
-                cpp_binaryop.binary_operation(
-                    lhs.view(),
-                    dereference(rhs.c_obj),
-                    op,
-                    output_type.c_obj
-                )
+            result = cpp_binaryop.binary_operation(
+                lhs.view(),
+                dereference(rhs.c_obj),
+                op,
+                output_type.c_obj
             )
     elif LeftBinaryOperand is Scalar and RightBinaryOperand is Column:
         with nogil:
-            result = move(
-                cpp_binaryop.binary_operation(
-                    dereference(lhs.c_obj),
-                    rhs.view(),
-                    op,
-                    output_type.c_obj
-                )
+            result = cpp_binaryop.binary_operation(
+                dereference(lhs.c_obj),
+                rhs.view(),
+                op,
+                output_type.c_obj
             )
     else:
         raise ValueError(f"Invalid arguments {lhs} and {rhs}")

python/pylibcudf/pylibcudf/column.pyx

@@ -138,7 +138,7 @@ cdef class Column:
 
         cdef size_type null_count = libcudf_col.get().null_count()
 
-        cdef column_contents contents = move(libcudf_col.get().release())
+        cdef column_contents contents = libcudf_col.get().release()
 
         # Note that when converting to cudf Column objects we'll need to pull
         # out the base object.
@@ -247,7 +247,7 @@ cdef class Column:
         cdef const scalar* c_scalar = slr.get()
         cdef unique_ptr[column] c_result
         with nogil:
-            c_result = move(make_column_from_scalar(dereference(c_scalar), size))
+            c_result = make_column_from_scalar(dereference(c_scalar), size)
         return Column.from_libcudf(move(c_result))
 
     @staticmethod
@@ -269,7 +269,7 @@ cdef class Column:
         cdef Scalar slr = Scalar.empty_like(like)
         cdef unique_ptr[column] c_result
         with nogil:
-            c_result = move(make_column_from_scalar(dereference(slr.get()), size))
+            c_result = make_column_from_scalar(dereference(slr.get()), size)
         return Column.from_libcudf(move(c_result))
 
     @staticmethod
@@ -373,7 +373,7 @@ cdef class Column:
         """Create a copy of the column."""
         cdef unique_ptr[column] c_result
         with nogil:
-            c_result = move(make_unique[column](self.view()))
+            c_result = make_unique[column](self.view())
         return Column.from_libcudf(move(c_result))
 
 

python/pylibcudf/pylibcudf/column_factories.pxd

@@ -1,7 +1,5 @@
 # Copyright (c) 2024, NVIDIA CORPORATION.
-from libcpp.memory cimport unique_ptr
-from libcpp.utility cimport move
-from pylibcudf.libcudf.types cimport mask_state, size_type
+from pylibcudf.libcudf.types cimport mask_state
 
 from .column cimport Column
 from .types cimport DataType, size_type, type_id

python/pylibcudf/pylibcudf/column_factories.pyx

@@ -39,29 +39,17 @@ cpdef Column make_empty_column(MakeEmptyColumnOperand type_or_id):
         if isinstance(type_or_id, TypeId):
             id = type_or_id
             with nogil:
-                result = move(
-                    cpp_make_empty_column(
-                        id
-                    )
-                )
+                result = cpp_make_empty_column(id)
         else:
             raise TypeError(
                 "Must pass a TypeId or DataType"
             )
     elif MakeEmptyColumnOperand is DataType:
         with nogil:
-            result = move(
-                cpp_make_empty_column(
-                    type_or_id.c_obj
-                )
-            )
+            result = cpp_make_empty_column(type_or_id.c_obj)
     elif MakeEmptyColumnOperand is type_id:
         with nogil:
-            result = move(
-                cpp_make_empty_column(
-                    type_or_id
-                )
-            )
+            result = cpp_make_empty_column(type_or_id)
     else:
         raise TypeError(
             "Must pass a TypeId or DataType"
@@ -92,12 +80,10 @@ cpdef Column make_numeric_column(
     else:
         raise TypeError("Invalid mask argument")
     with nogil:
-        result = move(
-            cpp_make_numeric_column(
-                type_.c_obj,
-                size,
-                state
-            )
+        result = cpp_make_numeric_column(
+            type_.c_obj,
+            size,
+            state
         )
 
     return Column.from_libcudf(move(result))
@@ -121,12 +107,10 @@ cpdef Column make_fixed_point_column(
     else:
         raise TypeError("Invalid mask argument")
     with nogil:
-        result = move(
-            cpp_make_fixed_point_column(
-                type_.c_obj,
-                size,
-                state
-            )
+        result = cpp_make_fixed_point_column(
+            type_.c_obj,
+            size,
+            state
         )
 
     return Column.from_libcudf(move(result))
@@ -151,12 +135,10 @@ cpdef Column make_timestamp_column(
     else:
         raise TypeError("Invalid mask argument")
     with nogil:
-        result = move(
-            cpp_make_timestamp_column(
-                type_.c_obj,
-                size,
-                state
-            )
+        result = cpp_make_timestamp_column(
+            type_.c_obj,
+            size,
+            state
         )
 
     return Column.from_libcudf(move(result))
@@ -181,12 +163,10 @@ cpdef Column make_duration_column(
     else:
         raise TypeError("Invalid mask argument")
     with nogil:
-        result = move(
-            cpp_make_duration_column(
-                type_.c_obj,
-                size,
-                state
-            )
+        result = cpp_make_duration_column(
+            type_.c_obj,
+            size,
+            state
         )
 
     return Column.from_libcudf(move(result))
@@ -211,12 +191,10 @@ cpdef Column make_fixed_width_column(
     else:
         raise TypeError("Invalid mask argument")
     with nogil:
-        result = move(
-            cpp_make_fixed_width_column(
-                type_.c_obj,
-                size,
-                state
-            )
+        result = cpp_make_fixed_width_column(
+            type_.c_obj,
+            size,
+            state
         )
 
     return Column.from_libcudf(move(result))

python/pylibcudf/pylibcudf/concatenate.pyx

@@ -40,14 +40,14 @@ cpdef concatenate(list objects):
             c_tables.push_back((<Table?>tbl).view())
 
         with nogil:
-            c_tbl_result = move(cpp_concatenate.concatenate(c_tables))
+            c_tbl_result = cpp_concatenate.concatenate(c_tables)
         return Table.from_libcudf(move(c_tbl_result))
     elif isinstance(objects[0], Column):
         for column in objects:
             c_columns.push_back((<Column?>column).view())
 
         with nogil:
-            c_col_result = move(cpp_concatenate.concatenate(c_columns))
+            c_col_result = cpp_concatenate.concatenate(c_columns)
         return Column.from_libcudf(move(c_col_result))
     else:
         raise ValueError("input must be a list of Columns or Tables")