scalar_device_view.cuh

Go to the documentation of this file.

/*
 * Copyright (c) 2019-2023, 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.
 */
#pragma once
 
#include <cudf/scalar/scalar.hpp>
#include <cudf/strings/string_view.hpp>
#include <cudf/types.hpp>
 
namespace cudf {
namespace detail {
class scalar_device_view_base {
 public:
  ~scalar_device_view_base() = default;
 
  [[nodiscard]] __host__ __device__ data_type type() const noexcept { return _type; }
 
  [[nodiscard]] __device__ bool is_valid() const noexcept { return *_is_valid; }
 
  __device__ void set_valid(bool is_valid) noexcept { *_is_valid = is_valid; }
 
 protected:
  data_type _type{type_id::EMPTY};  
  bool* _is_valid{};                
 
  scalar_device_view_base(data_type type, bool* is_valid) : _type(type), _is_valid(is_valid) {}
 
  scalar_device_view_base() = default;
};
 
class fixed_width_scalar_device_view_base : public detail::scalar_device_view_base {
 public:
  template <typename T>
  __device__ T& value() noexcept
  {
    return *data<T>();
  }
 
  template <typename T>
  __device__ T const& value() const noexcept
  {
    return *data<T>();
  }
 
  template <typename T>
  __device__ void set_value(T value)
  {
    *static_cast<T*>(_data) = value;
  }
 
  template <typename T>
  __device__ T* data() noexcept
  {
    return static_cast<T*>(_data);
  }
 
  template <typename T>
  __device__ T const* data() const noexcept
  {
    return static_cast<T const*>(_data);
  }
 
 protected:
  void* _data{};  
 
  fixed_width_scalar_device_view_base(data_type type, void* data, bool* is_valid)
    : detail::scalar_device_view_base(type, is_valid), _data(data)
  {
  }
};
 
template <typename T>
class fixed_width_scalar_device_view : public detail::fixed_width_scalar_device_view_base {
 public:
  using value_type = T;  
 
  __device__ T& value() noexcept { return fixed_width_scalar_device_view_base::value<T>(); }
 
  __device__ T const& value() const noexcept
  {
    return fixed_width_scalar_device_view_base::value<T>();
  }
 
  __device__ void set_value(T value) { fixed_width_scalar_device_view_base::set_value<T>(value); }
 
  __device__ T* data() noexcept { return fixed_width_scalar_device_view_base::data<T>(); }
 
  __device__ T const* data() const noexcept
  {
    return fixed_width_scalar_device_view_base::data<T>();
  }
 
 protected:
  fixed_width_scalar_device_view(data_type type, T* data, bool* is_valid)
    : detail::fixed_width_scalar_device_view_base(type, data, is_valid)
  {
  }
};
 
}  // namespace detail
 
template <typename T>
class numeric_scalar_device_view : public detail::fixed_width_scalar_device_view<T> {
 public:
  numeric_scalar_device_view(data_type type, T* data, bool* is_valid)
    : detail::fixed_width_scalar_device_view<T>(type, data, is_valid)
  {
  }
};
 
template <typename T>
class fixed_point_scalar_device_view : public detail::scalar_device_view_base {
 public:
  using rep_type = typename T::rep;  
 
  fixed_point_scalar_device_view(data_type type, rep_type* data, bool* is_valid)
    : detail::scalar_device_view_base(type, is_valid), _data(data)
  {
  }
 
  __device__ void set_value(rep_type value) { *_data = value; }
 
  __device__ rep_type const& rep() const noexcept { return *_data; }
 
 private:
  rep_type* _data{};
};
 
class string_scalar_device_view : public detail::scalar_device_view_base {
 public:
  using ValueType = cudf::string_view;  
 
  string_scalar_device_view(data_type type, char const* data, bool* is_valid, size_type size)
    : detail::scalar_device_view_base(type, is_valid), _data(data), _size(size)
  {
  }
 
  [[nodiscard]] __device__ ValueType value() const noexcept
  {
    return ValueType{this->data(), _size};
  }
 
  [[nodiscard]] __device__ char const* data() const noexcept
  {
    return static_cast<char const*>(_data);
  }
 
  [[nodiscard]] __device__ size_type size() const noexcept { return _size; }
 
 private:
  char const* _data{};  
  size_type _size;      
};
 
template <typename T>
class timestamp_scalar_device_view : public detail::fixed_width_scalar_device_view<T> {
 public:
  timestamp_scalar_device_view(data_type type, T* data, bool* is_valid)
    : detail::fixed_width_scalar_device_view<T>(type, data, is_valid)
  {
  }
};
 
template <typename T>
class duration_scalar_device_view : public detail::fixed_width_scalar_device_view<T> {
 public:
  duration_scalar_device_view(data_type type, T* data, bool* is_valid)
    : detail::fixed_width_scalar_device_view<T>(type, data, is_valid)
  {
  }
};
 
template <typename T>
auto get_scalar_device_view(numeric_scalar<T>& s)
{
  return numeric_scalar_device_view<T>(s.type(), s.data(), s.validity_data());
}
 
inline auto get_scalar_device_view(string_scalar& s)
{
  return string_scalar_device_view(s.type(), s.data(), s.validity_data(), s.size());
}
 
template <typename T>
auto get_scalar_device_view(timestamp_scalar<T>& s)
{
  return timestamp_scalar_device_view<T>(s.type(), s.data(), s.validity_data());
}
 
template <typename T>
auto get_scalar_device_view(duration_scalar<T>& s)
{
  return duration_scalar_device_view<T>(s.type(), s.data(), s.validity_data());
}
 
template <typename T>
auto get_scalar_device_view(fixed_point_scalar<T>& s)
{
  return fixed_point_scalar_device_view<T>(s.type(), s.data(), s.validity_data());
}
 
}  // namespace cudf