Add a public API for copying a table_view to device array

cpp/CMakeLists.txt

@@ -644,6 +644,7 @@ add_library(
   src/reshape/byte_cast.cu
   src/reshape/interleave_columns.cu
   src/reshape/tile.cu
+  src/reshape/table_to_array.cu
   src/rolling/detail/optimized_unbounded_window.cpp
   src/rolling/detail/rolling_collect_list.cu
   src/rolling/detail/rolling_fixed_window.cu

cpp/benchmarks/CMakeLists.txt

@@ -430,7 +430,7 @@ ConfigureNVBench(DECIMAL_NVBENCH decimal/convert_floating.cpp)
 # ##################################################################################################
 # * reshape benchmark
 # ---------------------------------------------------------------------------------
-ConfigureNVBench(RESHAPE_NVBENCH reshape/interleave.cpp)
+ConfigureNVBench(RESHAPE_NVBENCH reshape/interleave.cpp reshape/table_to_array.cpp)
 
 # ##################################################################################################
 # * rolling benchmark

cpp/benchmarks/reshape/table_to_array.cpp

@@ -0,0 +1,58 @@
+/*
+ * Copyright (c) 2025, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <benchmarks/common/generate_input.hpp>
+
+#include <cudf/reshape.hpp>
+#include <cudf/utilities/default_stream.hpp>
+
+#include <nvbench/nvbench.cuh>
+
+static void bench_table_to_array(nvbench::state& state)
+{
+  auto const num_rows = static_cast<cudf::size_type>(state.get_int64("num_rows"));
+  auto const num_cols = static_cast<cudf::size_type>(state.get_int64("columns"));
+
+  if (static_cast<std::size_t>(num_rows) * num_cols >=
+      static_cast<std::size_t>(std::numeric_limits<cudf::size_type>::max())) {
+    state.skip("Input size exceeds cudf::size_type limit");
+  }
+
+  data_profile profile =
+    data_profile_builder().distribution(cudf::type_id::INT32, distribution_id::UNIFORM, 0, 1000);
+  profile.set_null_probability(0.0);
+  std::vector<cudf::type_id> types(num_cols, cudf::type_id::INT32);
+  auto input_table = create_random_table(types, row_count{num_rows}, profile);
+
+  auto input_view = input_table->view();
+  auto stream     = cudf::get_default_stream();
+  auto dtype      = cudf::data_type{cudf::type_id::INT32};
+
+  rmm::device_buffer output(num_rows * num_cols * sizeof(int32_t), stream);
+
+  state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value()));
+  state.add_global_memory_reads<int32_t>(num_rows * num_cols);   // all bytes are read
+  state.add_global_memory_writes<int32_t>(num_rows * num_cols);  // all bytes are written
+
+  state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
+    cudf::table_to_array(input_view, output.data(), dtype, stream);
+  });
+}
+
+NVBENCH_BENCH(bench_table_to_array)
+  .set_name("table_to_array")
+  .add_int64_axis("num_rows", {32768, 262144, 2097152, 16777216})
+  .add_int64_axis("columns", {2, 10, 100});

cpp/include/cudf/detail/reshape.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -41,5 +41,13 @@ std::unique_ptr<column> interleave_columns(table_view const& input,
                                            rmm::cuda_stream_view,
                                            rmm::device_async_resource_ref mr);
 
+/**
+ * @copydoc cudf::table_to_array
+ */
+void table_to_array(table_view const& input,
+                    void* output,
+                    data_type output_dtype,
+                    rmm::cuda_stream_view stream);
+
 }  // namespace detail
 }  // namespace CUDF_EXPORT cudf

cpp/include/cudf/reshape.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -107,6 +107,26 @@ std::unique_ptr<column> byte_cast(
   rmm::cuda_stream_view stream      = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
+/**
+ * @brief Copies a table into a contiguous column-major device array.
+ *
+ * This function converts a table_view with columns of the same type
+ * into a 2D device array in column-major order. The output buffer must be
+ * preallocated and large enough to hold `num_rows * num_columns` values of `output_dtype`.
+ *
+ * @throws cudf::logic_error if column types do not match `output_dtype`
+ * @throws cudf::logic_error if `output_dtype` is not a fixed-width type
+ *
+ * @param input A table with fixed-width, non-nullable columns of the same type
+ * @param output Pointer to device memory sized to hold `num_rows * num_columns` values
+ * @param output_dtype The data type of the output array
+ * @param stream CUDA stream used for memory operations
+ */
+void table_to_array(table_view const& input,
+                    void* output,
+                    cudf::data_type output_dtype,
+                    rmm::cuda_stream_view stream = cudf::get_default_stream());
+
 /** @} */  // end of group
 
 }  // namespace CUDF_EXPORT cudf

cpp/src/reshape/table_to_array.cu

@@ -0,0 +1,146 @@
+/*
+ * Copyright (c) 2025, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/reshape.hpp>
+#include <cudf/reshape.hpp>
+#include <cudf/utilities/default_stream.hpp>
+#include <cudf/utilities/error.hpp>
+#include <cudf/utilities/type_dispatcher.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/device_uvector.hpp>
+
+#include <cub/device/device_memcpy.cuh>
+#include <cuda/functional>
+#include <cuda_runtime.h>
+#include <thrust/device_vector.h>
+#include <thrust/iterator/constant_iterator.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/transform_iterator.h>
+
+namespace cudf {
+namespace detail {
+namespace {
+
+template <typename T>
+void _table_to_array(table_view const& input, void* output, rmm::cuda_stream_view stream)
+{
+  auto const num_columns = input.num_columns();
+  auto const num_rows    = input.num_rows();
+  auto const item_size   = sizeof(T);
+  auto* base_ptr         = static_cast<cuda::std::byte*>(output);
+
+  CUDF_EXPECTS(num_columns > 0, "Must have at least one column.");
+  CUDF_EXPECTS(output != nullptr, "Output pointer cannot be null.");
+
+  rmm::device_uvector<void*> d_srcs(num_columns, stream);
+  rmm::device_uvector<void*> d_dsts(num_columns, stream);
+
+  std::vector<void const*> h_srcs(num_columns);
+  std::vector<void*> h_dsts(num_columns);
+
+  for (int i = 0; i < num_columns; ++i) {
+    auto const& col = input.column(i);
+    CUDF_EXPECTS(col.type() == input.column(0).type(), "All columns must have the same dtype");
+    CUDF_EXPECTS(col.null_count() == 0, "All columns must be non-nullable or contain no nulls");
+
+    h_srcs[i] = static_cast<void const*>(col.data<T>());
+    h_dsts[i] = static_cast<void*>(base_ptr + i * item_size * num_rows);
+  }
+
+  CUDF_CUDA_TRY(cudaMemcpyAsync(d_srcs.data(),
+                                h_srcs.data(),
+                                sizeof(void*) * num_columns,
+                                cudaMemcpyHostToDevice,
+                                stream.value()));
+  CUDF_CUDA_TRY(cudaMemcpyAsync(d_dsts.data(),
+                                h_dsts.data(),
+                                sizeof(void*) * num_columns,
+                                cudaMemcpyHostToDevice,
+                                stream.value()));
+
+  thrust::constant_iterator<size_t> sizes(static_cast<size_t>(item_size * num_rows));
+
+  void* d_temp_storage      = nullptr;
+  size_t temp_storage_bytes = 0;
+  cub::DeviceMemcpy::Batched(d_temp_storage,
+                             temp_storage_bytes,
+                             d_srcs.begin(),
+                             d_dsts.begin(),
+                             sizes,
+                             num_columns,
+                             stream.value());
+
+  rmm::device_buffer temp_storage(temp_storage_bytes, stream);
+
+  cub::DeviceMemcpy::Batched(temp_storage.data(),
+                             temp_storage_bytes,
+                             d_srcs.begin(),
+                             d_dsts.begin(),
+                             sizes,
+                             num_columns,
+                             stream.value());
+}
+
+struct TableToArrayDispatcher {
+  table_view const& input;
+  void* output;
+  rmm::cuda_stream_view stream;
+
+  template <typename T, CUDF_ENABLE_IF(is_fixed_width<T>() || is_fixed_point<T>())>
+  void operator()() const
+  {
+    if constexpr (is_fixed_point<T>()) {
+      using StorageType = cudf::device_storage_type_t<T>;
+      _table_to_array<StorageType>(input, output, stream);
+    } else {
+      _table_to_array<T>(input, output, stream);
+    }
+  }
+
+  template <typename T, CUDF_ENABLE_IF(!is_fixed_width<T>() && !is_fixed_point<T>())>
+  void operator()() const
+  {
+    CUDF_FAIL("Unsupported dtype");
+  }
+};
+
+}  // namespace
+
+void table_to_array(table_view const& input,
+                    void* output,
+                    data_type output_dtype,
+                    rmm::cuda_stream_view stream)
+{
+  CUDF_EXPECTS(output != nullptr, "Output pointer cannot be null.");
+  CUDF_EXPECTS(input.num_columns() > 0, "Input must have at least one column.");
+
+  cudf::type_dispatcher(output_dtype, TableToArrayDispatcher{input, output, stream});
+}
+
+}  // namespace detail
+
+void table_to_array(table_view const& input,
+                    void* output,
+                    data_type output_dtype,
+                    rmm::cuda_stream_view stream)
+{
+  CUDF_FUNC_RANGE();
+  cudf::detail::table_to_array(input, output, output_dtype, stream);
+}
+
+}  // namespace cudf

cpp/tests/CMakeLists.txt

@@ -520,7 +520,7 @@ ConfigureTest(
 # * reshape test ----------------------------------------------------------------------------------
 ConfigureTest(
   RESHAPE_TEST reshape/byte_cast_tests.cpp reshape/interleave_columns_tests.cpp
-  reshape/tile_tests.cpp
+  reshape/tile_tests.cpp reshape/table_to_array.cpp
 )
 
 # ##################################################################################################