column_device_view_base.cuh

Go to the documentation of this file.

/*
 * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: Apache-2.0
 */
#pragma once
 
#include <cudf/column/column_child_offsets.hpp>
#include <cudf/detail/offsets_iterator.cuh>
#include <cudf/fixed_point/fixed_point.hpp>
#include <cudf/strings/string_view.cuh>
#include <cudf/types.hpp>
#include <cudf/utilities/bit.hpp>
#include <cudf/utilities/traits.hpp>
 
#include <cuda/std/algorithm>
#include <cuda/std/optional>
#include <cuda/std/type_traits>
 
namespace CUDF_EXPORT cudf {
 
struct nullate {
  struct YES : cuda::std::true_type {};
  struct NO : cuda::std::false_type {};
  struct DYNAMIC {
    DYNAMIC() = delete;
    constexpr explicit DYNAMIC(bool b) noexcept : value{b} {}
    CUDF_HOST_DEVICE constexpr operator bool() const noexcept { return value; }
    bool value;  
  };
};
 
template <typename IndexType, typename KeyType>
  requires(is_index_type<IndexType>() && is_relationally_comparable<KeyType, KeyType>())
struct dictionary_element {
  using index_type = IndexType;  
  using key_type   = KeyType;    
 
  key_type key{};  
};
 
template <typename T>
inline constexpr bool is_dictionary_encoded = false;
 
template <typename IndexType, typename KeyType>
inline constexpr bool is_dictionary_encoded<dictionary_element<IndexType, KeyType>> = true;
 
namespace detail {
class alignas(16) column_device_view_base {
 public:
  static constexpr size_type offsets_column_index =
    cudf::offsets_column_index;  
 
  column_device_view_base()                               = delete;
  ~column_device_view_base()                              = default;
  column_device_view_base(column_device_view_base const&) = default;  
  column_device_view_base(column_device_view_base&&)      = default;  
  column_device_view_base& operator=(column_device_view_base const&) = default;
  column_device_view_base& operator=(column_device_view_base&&) = default;
 
  template <typename T = void,
            CUDF_ENABLE_IF(cuda::std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
  [[nodiscard]] CUDF_HOST_DEVICE T const* head() const noexcept
  {
    return static_cast<T const*>(_data);
  }
 
  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
  [[nodiscard]] CUDF_HOST_DEVICE T const* data() const noexcept
  {
    return head<T>() + _offset;
  }
 
  [[nodiscard]] CUDF_HOST_DEVICE size_type size() const noexcept { return _size; }
 
  [[nodiscard]] CUDF_HOST_DEVICE data_type type() const noexcept { return _type; }
 
  [[nodiscard]] CUDF_HOST_DEVICE bool nullable() const noexcept { return nullptr != _null_mask; }
 
  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type const* null_mask() const noexcept
  {
    return _null_mask;
  }
 
  [[nodiscard]] CUDF_HOST_DEVICE size_type offset() const noexcept { return _offset; }
 
  [[nodiscard]] __device__ bool is_valid(size_type element_index) const noexcept
  {
    return not nullable() or is_valid_nocheck(element_index);
  }
 
  [[nodiscard]] __device__ bool is_valid_nocheck(size_type element_index) const noexcept
  {
    return bit_is_set(_null_mask, offset() + element_index);
  }
 
  [[nodiscard]] __device__ bool is_null(size_type element_index) const noexcept
  {
    return not is_valid(element_index);
  }
 
  [[nodiscard]] __device__ bool is_null_nocheck(size_type element_index) const noexcept
  {
    return not is_valid_nocheck(element_index);
  }
 
  [[nodiscard]] __device__ bitmask_type get_mask_word(size_type word_index) const noexcept
  {
    return null_mask()[word_index];
  }
 
 protected:
  data_type _type{type_id::EMPTY};   
  cudf::size_type _size{};           
  void const* _data{};               
  size_type _null_count{};           
  bitmask_type const* _null_mask{};  
  size_type _offset{};               
  void* _children{};                 
  size_type _num_children{};         
 
  CUDF_HOST_DEVICE column_device_view_base(data_type type,
                                           size_type size,
                                           void const* data,
                                           size_type null_count,
                                           bitmask_type const* null_mask,
                                           size_type offset,
                                           void* children,
                                           size_type num_children)
    : _type{type},
      _size{size},
      _data{data},
      _null_count{null_count},
      _null_mask{null_mask},
      _offset{offset},
      _children{children},
      _num_children{num_children}
  {
  }
 
  template <typename C, typename T, typename = void>
  struct has_element_accessor_impl : cuda::std::false_type {};
 
  template <typename C, typename T>
  struct has_element_accessor_impl<
    C,
    T,
    void_t<decltype(cuda::std::declval<C>().template element<T>(cuda::std::declval<size_type>()))>>
    : cuda::std::true_type {};
};
// @cond
// Forward declaration
template <typename T>
struct value_accessor;
template <typename T, typename Nullate>
struct optional_accessor;
template <typename T, bool has_nulls>
struct pair_accessor;
template <typename T, bool has_nulls>
struct pair_rep_accessor;
template <typename T>
struct mutable_value_accessor;
// @endcond
}  // namespace detail
 
class alignas(16) column_device_view_core : public detail::column_device_view_base {
 public:
  static constexpr bool is_mutable =
    false;  
 
  column_device_view_core()                               = delete;
  ~column_device_view_core()                              = default;
  column_device_view_core(column_device_view_core const&) = default;  
  column_device_view_core(column_device_view_core&&)      = default;  
  column_device_view_core& operator=(column_device_view_core const&) = default;
  column_device_view_core& operator=(column_device_view_core&&) = default;
 
  column_device_view_core(column_view column, void* h_ptr, void* d_ptr);
 
  [[nodiscard]] CUDF_HOST_DEVICE column_device_view_core slice(size_type offset,
                                                               size_type size) const noexcept
  {
    return column_device_view_core{this->type(),
                                   size,
                                   this->head(),
                                   this->null_count(),
                                   this->null_mask(),
                                   this->offset() + offset,
                                   static_cast<column_device_view_core*>(_children),
                                   this->num_child_columns()};
  }
 
  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
  {
    return data<T>()[element_index];
  }
 
  template <typename T, CUDF_ENABLE_IF(cuda::std::is_same_v<T, string_view>)>
  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
  {
    size_type index       = element_index + offset();  // account for this view's _offset
    char const* d_strings = static_cast<char const*>(_data);
    auto const offsets    = child(offsets_column_index);
    auto const itr        = cudf::detail::input_offsetalator(offsets.head(), offsets.type());
    auto const offset     = itr[index];
    return string_view{d_strings + offset, static_cast<cudf::size_type>(itr[index + 1] - offset)};
  }
 
 public:
  template <typename T, CUDF_ENABLE_IF(cudf::is_fixed_point<T>())>
  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
  {
    using namespace numeric;
    using rep        = typename T::rep;
    auto const scale = scale_type{_type.scale()};
    return T{scaled_integer<rep>{data<rep>()[element_index], scale}};
  }
 
  template <typename T, CUDF_ENABLE_IF(is_dictionary_encoded<T>)>
  [[nodiscard]] __device__ decltype(auto) element(size_type element_index) const noexcept
  {
    auto const& indices = child(dictionary_indices_column_index);
    auto const& keys    = child(dictionary_keys_column_index);
    auto const index    = indices.template element<typename T::index_type>(
      element_index + offset());  // account for this view's _offset
    return keys.template element<typename T::key_type>(index);
  }
 
  template <typename T>
  [[nodiscard]] __device__ cuda::std::optional<T> nullable_element(
    size_type element_index) const noexcept
  {
    if (is_null(element_index)) { return cuda::std::nullopt; }
    return element<T>(element_index);
  }
 
  [[nodiscard]] __device__ column_device_view_core child(size_type child_index) const noexcept
  {
    return static_cast<column_device_view_core*>(_children)[child_index];
  }
 
  [[nodiscard]] CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
  {
    return _num_children;
  }
 
  [[nodiscard]] CUDF_HOST_DEVICE size_type null_count() const noexcept { return _null_count; }
 
 protected:
  CUDF_HOST_DEVICE column_device_view_core(data_type type,
                                           size_type size,
                                           void const* data,
                                           size_type null_count,
                                           bitmask_type const* null_mask,
                                           size_type offset,
                                           column_device_view_core* children,
                                           size_type num_children)
    : column_device_view_base(
        type, size, data, null_count, null_mask, offset, children, num_children)
  {
  }
};
 
class alignas(16) mutable_column_device_view_core : public detail::column_device_view_base {
 public:
  static constexpr bool is_mutable =
    true;  
 
  mutable_column_device_view_core()  = delete;
  ~mutable_column_device_view_core() = default;
  mutable_column_device_view_core(mutable_column_device_view_core const&) =
    default;  
  mutable_column_device_view_core(mutable_column_device_view_core&&) =
    default;  
  mutable_column_device_view_core& operator=(mutable_column_device_view_core const&) = default;
  mutable_column_device_view_core& operator=(mutable_column_device_view_core&&) = default;
 
  template <typename T = void,
            CUDF_ENABLE_IF(cuda::std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
  [[nodiscard]] CUDF_HOST_DEVICE T* head() const noexcept
  {
    return const_cast<T*>(detail::column_device_view_base::head<T>());
  }
 
  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
  [[nodiscard]] CUDF_HOST_DEVICE T* data() const noexcept
  {
    return const_cast<T*>(detail::column_device_view_base::data<T>());
  }
 
  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
  [[nodiscard]] __device__ T& element(size_type element_index) const noexcept
  {
    return data<T>()[element_index];
  }
 
  template <typename T, CUDF_ENABLE_IF(cuda::std::is_same_v<T, string_view>)>
  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
  {
    size_type index       = element_index + offset();  // account for this view's _offset
    char const* d_strings = static_cast<char const*>(_data);
    auto const offsets    = child(offsets_column_index);
    auto const itr        = cudf::detail::input_offsetalator(offsets.head(), offsets.type());
    auto const offset     = itr[index];
    return string_view{d_strings + offset, static_cast<cudf::size_type>(itr[index + 1] - offset)};
  }
 
  template <typename T, CUDF_ENABLE_IF(cudf::is_fixed_point<T>())>
  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
  {
    using namespace numeric;
    using rep        = typename T::rep;
    auto const scale = scale_type{_type.scale()};
    return T{scaled_integer<rep>{data<rep>()[element_index], scale}};
  }
 
  template <typename T>
  [[nodiscard]] __device__ cuda::std::optional<T> nullable_element(
    size_type element_index) const noexcept
  {
    if (is_null(element_index)) { return cuda::std::nullopt; }
    return element<T>(element_index);
  }
 
  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
  __device__ void assign(size_type element_index, T value) const noexcept
  {
    data<T>()[element_index] = value;
  }
 
  template <typename T, CUDF_ENABLE_IF(is_fixed_point<T>())>
  __device__ void assign(size_type element_index, T value) const noexcept
  {
    // consider asserting that the scale matches
    using namespace numeric;
    using rep                  = typename T::rep;
    data<rep>()[element_index] = value.value();
  }
 
  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type* null_mask() const noexcept
  {
    return const_cast<bitmask_type*>(detail::column_device_view_base::null_mask());
  }
 
  [[nodiscard]] __device__ mutable_column_device_view_core
  child(size_type child_index) const noexcept
  {
    return static_cast<mutable_column_device_view_core*>(_children)[child_index];
  }
 
#ifdef __CUDACC__  // because set_bit in bit.hpp is wrapped with __CUDACC__
  __device__ void set_valid(size_type element_index) const noexcept
  {
    return set_bit(null_mask(), element_index);
  }
 
  __device__ void set_null(size_type element_index) const noexcept
  {
    return clear_bit(null_mask(), element_index);
  }
 
#endif
 
  __device__ void set_mask_word(size_type word_index, bitmask_type new_word) const noexcept
  {
    null_mask()[word_index] = new_word;
  }
 
 protected:
  CUDF_HOST_DEVICE mutable_column_device_view_core(data_type type,
                                                   size_type size,
                                                   void const* data,
                                                   bitmask_type const* null_mask,
                                                   size_type offset,
                                                   mutable_column_device_view_core* children,
                                                   size_type num_children)
    : column_device_view_base(type,
                              size,
                              data,
                              0,  // unused
                              null_mask,
                              offset,
                              children,
                              num_children)
  {
  }
};
 
}  // namespace CUDF_EXPORT cudf