[POC] Add BINARY type to cuDF supporting 32-bit and 64-bit offsets

cpp/CMakeLists.txt

@@ -581,6 +581,10 @@ add_library(
   src/ast/expressions.cpp
   src/ast/jit/expressions.cpp
   src/ast/operators.cpp
+  src/binary/binary_column_factories.cpp
+  src/binary/binary_column_view.cpp
+  src/binary/copying.cu
+  src/binary/from_views.cu
   src/binaryop/binaryop.cpp
   src/binaryop/compiled/ATan2.cu
   src/binaryop/compiled/Add.cu

cpp/include/cudf/binary/binary_column_factories.hpp

@@ -0,0 +1,69 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/binary/binary_view.hpp>
+#include <cudf/column/column.hpp>
+#include <cudf/utilities/default_stream.hpp>
+#include <cudf/utilities/export.hpp>
+#include <cudf/utilities/memory_resource.hpp>
+#include <cudf/utilities/span.hpp>
+
+#include <rmm/device_buffer.hpp>
+
+#include <memory>
+
+/**
+ * @file
+ * @brief Factory functions for `BINARY` columns.
+ */
+
+namespace CUDF_EXPORT cudf {
+
+/**
+ * @brief Constructs a `BINARY` column by copying device-resident binary views.
+ *
+ * A view whose data pointer equals `null_placeholder.data()` becomes null.
+ *
+ * @param binary_views Device span of binary values
+ * @param null_placeholder View whose data pointer marks null values
+ * @param stream CUDA stream used for device operations
+ * @param mr Device memory resource used for output allocations
+ * @return Newly constructed `BINARY` column
+ */
+std::unique_ptr<column> make_binary_column(
+  device_span<binary_view const> binary_views,
+  binary_view null_placeholder,
+  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+
+/**
+ * @brief Constructs a `BINARY` column from offsets and a contiguous byte buffer.
+ *
+ * `offsets_column` must contain either `INT32` or `INT64` values and must not
+ * contain nulls. Nonnegative, monotonically nondecreasing offsets whose final
+ * value does not exceed `bytes_buffer.size()` are a caller precondition.
+ *
+ * Empty input is canonicalized to an empty `BINARY` column with no children.
+ *
+ * @param num_rows Number of binary values represented by the column
+ * @param offsets_column Offsets with `num_rows + 1` elements
+ * @param bytes_buffer Contiguous payload bytes
+ * @param null_count Number of null rows
+ * @param null_mask Row validity mask
+ * @return Newly constructed `BINARY` column
+ */
+std::unique_ptr<column> make_binary_column(size_type num_rows,
+                                           std::unique_ptr<column> offsets_column,
+                                           rmm::device_buffer&& bytes_buffer,
+                                           size_type null_count,
+                                           rmm::device_buffer&& null_mask);
+
+/**
+ * @brief Creates an empty `BINARY` column.
+ */
+std::unique_ptr<column> make_empty_binary_column();
+
+}  // namespace CUDF_EXPORT cudf

cpp/include/cudf/binary/binary_column_view.hpp

@@ -0,0 +1,57 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/column/column_view.hpp>
+#include <cudf/utilities/export.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+
+#include <cstdint>
+
+/**
+ * @file
+ * @brief Class definition for `cudf::binary_column_view`.
+ */
+
+namespace CUDF_EXPORT cudf {
+
+/**
+ * @brief A wrapper providing access to a `BINARY` column's offsets and payload.
+ */
+class binary_column_view : private column_view {
+ public:
+  static constexpr size_type offsets_column_index{0};
+
+  explicit binary_column_view(column_view binary_column);
+  binary_column_view()                          = default;
+  binary_column_view(binary_column_view const&) = default;
+  binary_column_view(binary_column_view&&)      = default;
+  ~binary_column_view() override                = default;
+  binary_column_view& operator=(binary_column_view const&) = default;
+  binary_column_view& operator=(binary_column_view&&) = default;
+
+  using column_view::has_nulls;
+  using column_view::is_empty;
+  using column_view::null_count;
+  using column_view::null_mask;
+  using column_view::offset;
+  using column_view::size;
+
+  [[nodiscard]] column_view parent() const;
+  [[nodiscard]] column_view offsets() const;
+
+  /**
+   * @brief Returns the total number of bytes in the underlying payload buffer.
+   *
+   * This reports the unsliced parent's payload size.
+   */
+  [[nodiscard]] int64_t bytes_size(rmm::cuda_stream_view stream) const;
+
+  [[nodiscard]] uint8_t const* bytes_begin() const noexcept;
+  [[nodiscard]] uint8_t const* bytes_end(rmm::cuda_stream_view stream) const;
+};
+
+}  // namespace CUDF_EXPORT cudf

cpp/include/cudf/binary/binary_view.hpp

@@ -0,0 +1,123 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/types.hpp>
+#include <cudf/utilities/export.hpp>
+
+#include <cuda/std/algorithm>
+
+#include <cstdint>
+
+/**
+ * @file
+ * @brief Class definition for `cudf::binary_view`.
+ */
+
+namespace CUDF_EXPORT cudf {
+
+/**
+ * @brief A non-owning, immutable view of a variable-length sequence of bytes.
+ *
+ * The caller must maintain the device memory for the lifetime of this object.
+ * Binary values have no character encoding; comparisons use unsigned-byte
+ * lexicographic ordering.
+ */
+class binary_view {
+ public:
+  using value_type     = uint8_t;
+  using const_iterator = value_type const*;
+
+  CUDF_HOST_DEVICE constexpr binary_view() = default;
+
+  /**
+   * @brief Constructs a view over `size` bytes beginning at `data`.
+   */
+  CUDF_HOST_DEVICE constexpr binary_view(value_type const* data, size_type size)
+    : _data{data}, _size{size}
+  {
+  }
+
+  /**
+   * @brief Returns the number of bytes in this value.
+   */
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr size_type size_bytes() const noexcept { return _size; }
+
+  /**
+   * @brief Returns a pointer to the first byte.
+   */
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr value_type const* data() const noexcept
+  {
+    return _data;
+  }
+
+  /**
+   * @brief Returns whether this value contains no bytes.
+   */
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr bool empty() const noexcept { return _size == 0; }
+
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr const_iterator begin() const noexcept { return _data; }
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr const_iterator end() const noexcept
+  {
+    return _data + _size;
+  }
+
+  /**
+   * @brief Returns the byte at `index`.
+   */
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr value_type operator[](size_type index) const noexcept
+  {
+    return _data[index];
+  }
+
+  /**
+   * @brief Compares values using unsigned-byte lexicographic ordering.
+   */
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr int compare(binary_view rhs) const noexcept
+  {
+    auto const count = cuda::std::min(_size, rhs._size);
+    for (size_type i = 0; i < count; ++i) {
+      if (_data[i] < rhs._data[i]) { return -1; }
+      if (_data[i] > rhs._data[i]) { return 1; }
+    }
+    return (_size > rhs._size) - (_size < rhs._size);
+  }
+
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator==(binary_view rhs) const noexcept
+  {
+    return compare(rhs) == 0;
+  }
+
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator!=(binary_view rhs) const noexcept
+  {
+    return not(*this == rhs);
+  }
+
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator<(binary_view rhs) const noexcept
+  {
+    return compare(rhs) < 0;
+  }
+
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator>(binary_view rhs) const noexcept
+  {
+    return compare(rhs) > 0;
+  }
+
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator<=(binary_view rhs) const noexcept
+  {
+    return compare(rhs) <= 0;
+  }
+
+  CUDF_HOST_DEVICE [[nodiscard]] constexpr bool operator>=(binary_view rhs) const noexcept
+  {
+    return compare(rhs) >= 0;
+  }
+
+ private:
+  value_type const* _data{};
+  size_type _size{};
+};
+
+}  // namespace CUDF_EXPORT cudf

cpp/include/cudf/binary/detail/binary_column_factories.cuh

@@ -0,0 +1,66 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/binary/binary_column_factories.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/valid_if.cuh>
+#include <cudf/strings/detail/strings_children.cuh>
+
+#include <cuda/functional>
+#include <cuda/std/iterator>
+#include <cuda/std/utility>
+#include <thrust/iterator/transform_iterator.h>
+
+namespace cudf::binary::detail {
+
+using binary_index_pair = cuda::std::pair<uint8_t const*, size_type>;
+
+/**
+ * @brief Creates a BINARY column from device-accessible pointer/size pairs.
+ *
+ * A null pointer denotes a null row. The bytes referenced by non-null pointers
+ * are copied into the output payload.
+ */
+template <typename IndexPairIterator>
+std::unique_ptr<column> make_binary_column(IndexPairIterator begin,
+                                           IndexPairIterator end,
+                                           rmm::cuda_stream_view stream,
+                                           rmm::device_async_resource_ref mr)
+{
+  CUDF_FUNC_RANGE();
+  auto const row_count = static_cast<size_type>(cuda::std::distance(begin, end));
+  if (row_count == 0) { return make_empty_binary_column(); }
+
+  auto sizes = thrust::make_transform_iterator(
+    begin,
+    cuda::proclaim_return_type<size_type>([] __device__(binary_index_pair item) {
+      return item.first == nullptr ? size_type{0} : item.second;
+    }));
+  auto [offsets, bytes] = cudf::strings::detail::make_offsets_child_column(
+    sizes, sizes + row_count, stream, mr);
+
+  auto const validator = [] __device__(binary_index_pair item) { return item.first != nullptr; };
+  auto [null_mask, null_count] = cudf::detail::valid_if(begin, end, validator, stream, mr);
+  if (null_count == 0) { null_mask = rmm::device_buffer{0, stream, mr}; }
+
+  auto char_pairs = thrust::make_transform_iterator(
+    begin,
+    cuda::proclaim_return_type<cuda::std::pair<char const*, size_type>>(
+      [] __device__(binary_index_pair item) {
+        return cuda::std::pair<char const*, size_type>{
+          reinterpret_cast<char const*>(item.first), item.second};
+      }));
+  auto payload = cudf::strings::detail::make_chars_buffer(
+    offsets->view(), bytes, char_pairs, row_count, stream, mr);
+
+  return cudf::make_binary_column(row_count,
+                                  std::move(offsets),
+                                  payload.release(),
+                                  null_count,
+                                  std::move(null_mask));
+}
+
+}  // namespace cudf::binary::detail

cpp/include/cudf/binary/detail/copying.hpp

@@ -0,0 +1,31 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/binary/binary_column_view.hpp>
+#include <cudf/column/column.hpp>
+#include <cudf/utilities/export.hpp>
+#include <cudf/utilities/memory_resource.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+
+#include <memory>
+
+namespace cudf::binary::detail {
+
+/**
+ * @brief Copies rows `[start, end)` from a BINARY column into a new owning column.
+ *
+ * The output offsets are normalized to begin at zero and retain the input
+ * offsets width.
+ */
+CUDF_EXPORT std::unique_ptr<column> copy_slice(
+  binary_column_view const& input,
+  size_type start,
+  size_type end,
+  rmm::cuda_stream_view stream,
+  rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+
+}  // namespace cudf::binary::detail

cpp/include/cudf/column/column_device_view.cuh

@@ -4,6 +4,7 @@
  */
 #pragma once
 
+#include <cudf/binary/binary_view.hpp>
 #include <cudf/column/column_device_view_base.cuh>
 #include <cudf/column/column_view.hpp>
 #include <cudf/detail/utilities/alignment.hpp>
@@ -132,7 +133,9 @@ class alignas(16) column_device_view : public column_device_view_core {
    * @param element_index Position of the desired string element
    * @return string_view instance representing this element at this index
    */
-  template <typename T, CUDF_ENABLE_IF(cuda::std::is_same_v<T, string_view>)>
+  template <typename T,
+            CUDF_ENABLE_IF(cuda::std::is_same_v<T, string_view> or
+                           cuda::std::is_same_v<T, binary_view>)>
   [[nodiscard]] __device__ T element(size_type element_index) const noexcept
   {
     return base::element<T>(element_index);