type_dispatcher.hpp

Go to the documentation of this file.

/*
 * Copyright (c) 2019-2022, 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/detail/utilities/assert.cuh>
#include <cudf/fixed_point/fixed_point.hpp>
#include <cudf/types.hpp>
#include <cudf/utilities/error.hpp>
#include <cudf/wrappers/dictionary.hpp>
#include <cudf/wrappers/durations.hpp>
#include <cudf/wrappers/timestamps.hpp>
 
#include <string>
 
namespace cudf {
template <typename T>
inline constexpr type_id type_to_id()
{
  return type_id::EMPTY;
};
 
struct type_to_name {
  template <typename T>
  inline std::string operator()()
  {
    return "void";
  }
};
 
template <cudf::type_id t>
struct id_to_type_impl {
  using type = void;
};
template <cudf::type_id Id>
using id_to_type = typename id_to_type_impl<Id>::type;
 
// clang-format off
template <typename T>
using device_storage_type_t =
  std::conditional_t<std::is_same_v<numeric::decimal32,  T>, int32_t,
  std::conditional_t<std::is_same_v<numeric::decimal64,  T>, int64_t,
  std::conditional_t<std::is_same_v<numeric::decimal128, T>, __int128_t, T>>>;
// clang-format on
 
template <typename T>
constexpr bool type_id_matches_device_storage_type(type_id id)
{
  return (id == type_id::DECIMAL32 && std::is_same_v<T, int32_t>) ||
         (id == type_id::DECIMAL64 && std::is_same_v<T, int64_t>) ||
         (id == type_id::DECIMAL128 && std::is_same_v<T, __int128_t>) || id == type_to_id<T>();
}
 
constexpr bool is_fixed_point(cudf::type_id id)
{
  return id == type_id::DECIMAL32 or  //
         id == type_id::DECIMAL64 or  //
         id == type_id::DECIMAL128;
}
 
#ifndef CUDF_TYPE_MAPPING
#define CUDF_TYPE_MAPPING(Type, Id)                   \
  template <>                                         \
  constexpr inline type_id type_to_id<Type>()         \
  {                                                   \
    return Id;                                        \
  }                                                   \
  template <>                                         \
  inline std::string type_to_name::operator()<Type>() \
  {                                                   \
    return CUDF_STRINGIFY(Type);                      \
  }                                                   \
  template <>                                         \
  struct id_to_type_impl<Id> {                        \
    using type = Type;                                \
  };
#endif
 
// Defines all of the mappings between C++ types and their corresponding `cudf::type_id` values.
CUDF_TYPE_MAPPING(bool, type_id::BOOL8)
CUDF_TYPE_MAPPING(int8_t, type_id::INT8)
CUDF_TYPE_MAPPING(int16_t, type_id::INT16)
CUDF_TYPE_MAPPING(int32_t, type_id::INT32)
CUDF_TYPE_MAPPING(int64_t, type_id::INT64)
CUDF_TYPE_MAPPING(uint8_t, type_id::UINT8)
CUDF_TYPE_MAPPING(uint16_t, type_id::UINT16)
CUDF_TYPE_MAPPING(uint32_t, type_id::UINT32)
CUDF_TYPE_MAPPING(uint64_t, type_id::UINT64)
CUDF_TYPE_MAPPING(float, type_id::FLOAT32)
CUDF_TYPE_MAPPING(double, type_id::FLOAT64)
CUDF_TYPE_MAPPING(cudf::string_view, type_id::STRING)
CUDF_TYPE_MAPPING(cudf::timestamp_D, type_id::TIMESTAMP_DAYS)
CUDF_TYPE_MAPPING(cudf::timestamp_s, type_id::TIMESTAMP_SECONDS)
CUDF_TYPE_MAPPING(cudf::timestamp_ms, type_id::TIMESTAMP_MILLISECONDS)
CUDF_TYPE_MAPPING(cudf::timestamp_us, type_id::TIMESTAMP_MICROSECONDS)
CUDF_TYPE_MAPPING(cudf::timestamp_ns, type_id::TIMESTAMP_NANOSECONDS)
CUDF_TYPE_MAPPING(cudf::duration_D, type_id::DURATION_DAYS)
CUDF_TYPE_MAPPING(cudf::duration_s, type_id::DURATION_SECONDS)
CUDF_TYPE_MAPPING(cudf::duration_ms, type_id::DURATION_MILLISECONDS)
CUDF_TYPE_MAPPING(cudf::duration_us, type_id::DURATION_MICROSECONDS)
CUDF_TYPE_MAPPING(cudf::duration_ns, type_id::DURATION_NANOSECONDS)
CUDF_TYPE_MAPPING(dictionary32, type_id::DICTIONARY32)
CUDF_TYPE_MAPPING(cudf::list_view, type_id::LIST)
CUDF_TYPE_MAPPING(numeric::decimal32, type_id::DECIMAL32)
CUDF_TYPE_MAPPING(numeric::decimal64, type_id::DECIMAL64)
CUDF_TYPE_MAPPING(numeric::decimal128, type_id::DECIMAL128)
CUDF_TYPE_MAPPING(cudf::struct_view, type_id::STRUCT)
 
 
template <cudf::type_id Id>
struct dispatch_storage_type {
  using type = device_storage_type_t<typename id_to_type_impl<Id>::type>;
};
 
template <typename T>
struct type_to_scalar_type_impl {
  using ScalarType = cudf::scalar;
};
 
#ifndef MAP_NUMERIC_SCALAR
#define MAP_NUMERIC_SCALAR(Type)                                     \
  template <>                                                        \
  struct type_to_scalar_type_impl<Type> {                            \
    using ScalarType       = cudf::numeric_scalar<Type>;             \
    using ScalarDeviceType = cudf::numeric_scalar_device_view<Type>; \
  };
#endif
 
MAP_NUMERIC_SCALAR(int8_t)
MAP_NUMERIC_SCALAR(int16_t)
MAP_NUMERIC_SCALAR(int32_t)
MAP_NUMERIC_SCALAR(int64_t)
MAP_NUMERIC_SCALAR(__int128_t)
MAP_NUMERIC_SCALAR(uint8_t)
MAP_NUMERIC_SCALAR(uint16_t)
MAP_NUMERIC_SCALAR(uint32_t)
MAP_NUMERIC_SCALAR(uint64_t)
MAP_NUMERIC_SCALAR(float)
MAP_NUMERIC_SCALAR(double)
MAP_NUMERIC_SCALAR(bool)
 
template <>
struct type_to_scalar_type_impl<std::string> {
  using ScalarType       = cudf::string_scalar;
  using ScalarDeviceType = cudf::string_scalar_device_view;
};
 
template <>
struct type_to_scalar_type_impl<cudf::string_view> {
  using ScalarType       = cudf::string_scalar;
  using ScalarDeviceType = cudf::string_scalar_device_view;
};
 
template <>
struct type_to_scalar_type_impl<numeric::decimal32> {
  using ScalarType       = cudf::fixed_point_scalar<numeric::decimal32>;
  using ScalarDeviceType = cudf::fixed_point_scalar_device_view<numeric::decimal32>;
};
 
template <>
struct type_to_scalar_type_impl<numeric::decimal64> {
  using ScalarType       = cudf::fixed_point_scalar<numeric::decimal64>;
  using ScalarDeviceType = cudf::fixed_point_scalar_device_view<numeric::decimal64>;
};
 
template <>
struct type_to_scalar_type_impl<numeric::decimal128> {
  using ScalarType       = cudf::fixed_point_scalar<numeric::decimal128>;
  using ScalarDeviceType = cudf::fixed_point_scalar_device_view<numeric::decimal128>;
};
 
template <>  // TODO: this is a temporary solution for make_pair_iterator
struct type_to_scalar_type_impl<cudf::dictionary32> {
  using ScalarType       = cudf::numeric_scalar<int32_t>;
  using ScalarDeviceType = cudf::numeric_scalar_device_view<int32_t>;
};
 
template <>  // TODO: this is to get compilation working. list scalars will be implemented at a
             // later time.
struct type_to_scalar_type_impl<cudf::list_view> {
  using ScalarType = cudf::list_scalar;
  // using ScalarDeviceType = cudf::list_scalar_device_view;
};
 
template <>  // TODO: Ditto, likewise.
struct type_to_scalar_type_impl<cudf::struct_view> {
  using ScalarType = cudf::struct_scalar;
  // using ScalarDeviceType = cudf::struct_scalar_device_view; // CALEB: TODO!
};
 
#ifndef MAP_TIMESTAMP_SCALAR
#define MAP_TIMESTAMP_SCALAR(Type)                                     \
  template <>                                                          \
  struct type_to_scalar_type_impl<Type> {                              \
    using ScalarType       = cudf::timestamp_scalar<Type>;             \
    using ScalarDeviceType = cudf::timestamp_scalar_device_view<Type>; \
  };
#endif
 
MAP_TIMESTAMP_SCALAR(timestamp_D)
MAP_TIMESTAMP_SCALAR(timestamp_s)
MAP_TIMESTAMP_SCALAR(timestamp_ms)
MAP_TIMESTAMP_SCALAR(timestamp_us)
MAP_TIMESTAMP_SCALAR(timestamp_ns)
 
 
#ifndef MAP_DURATION_SCALAR
#define MAP_DURATION_SCALAR(Type)                                     \
  template <>                                                         \
  struct type_to_scalar_type_impl<Type> {                             \
    using ScalarType       = cudf::duration_scalar<Type>;             \
    using ScalarDeviceType = cudf::duration_scalar_device_view<Type>; \
  };
#endif
 
MAP_DURATION_SCALAR(duration_D)
MAP_DURATION_SCALAR(duration_s)
MAP_DURATION_SCALAR(duration_ms)
MAP_DURATION_SCALAR(duration_us)
MAP_DURATION_SCALAR(duration_ns)
 
 
template <typename T>
using scalar_type_t = typename type_to_scalar_type_impl<T>::ScalarType;
 
template <typename T>
using scalar_device_type_t = typename type_to_scalar_type_impl<T>::ScalarDeviceType;
 
// This pragma disables a compiler warning that complains about the valid usage
// of calling a __host__ functor from this function which is __host__ __device__
#pragma nv_exec_check_disable
template <template <cudf::type_id> typename IdTypeMap = id_to_type_impl,
          typename Functor,
          typename... Ts>
CUDF_HOST_DEVICE __forceinline__ constexpr decltype(auto) type_dispatcher(cudf::data_type dtype,
                                                                          Functor f,
                                                                          Ts&&... args)
{
  switch (dtype.id()) {
    case type_id::BOOL8:
      return f.template operator()<typename IdTypeMap<type_id::BOOL8>::type>(
        std::forward<Ts>(args)...);
    case type_id::INT8:
      return f.template operator()<typename IdTypeMap<type_id::INT8>::type>(
        std::forward<Ts>(args)...);
    case type_id::INT16:
      return f.template operator()<typename IdTypeMap<type_id::INT16>::type>(
        std::forward<Ts>(args)...);
    case type_id::INT32:
      return f.template operator()<typename IdTypeMap<type_id::INT32>::type>(
        std::forward<Ts>(args)...);
    case type_id::INT64:
      return f.template operator()<typename IdTypeMap<type_id::INT64>::type>(
        std::forward<Ts>(args)...);
    case type_id::UINT8:
      return f.template operator()<typename IdTypeMap<type_id::UINT8>::type>(
        std::forward<Ts>(args)...);
    case type_id::UINT16:
      return f.template operator()<typename IdTypeMap<type_id::UINT16>::type>(
        std::forward<Ts>(args)...);
    case type_id::UINT32:
      return f.template operator()<typename IdTypeMap<type_id::UINT32>::type>(
        std::forward<Ts>(args)...);
    case type_id::UINT64:
      return f.template operator()<typename IdTypeMap<type_id::UINT64>::type>(
        std::forward<Ts>(args)...);
    case type_id::FLOAT32:
      return f.template operator()<typename IdTypeMap<type_id::FLOAT32>::type>(
        std::forward<Ts>(args)...);
    case type_id::FLOAT64:
      return f.template operator()<typename IdTypeMap<type_id::FLOAT64>::type>(
        std::forward<Ts>(args)...);
    case type_id::STRING:
      return f.template operator()<typename IdTypeMap<type_id::STRING>::type>(
        std::forward<Ts>(args)...);
    case type_id::TIMESTAMP_DAYS:
      return f.template operator()<typename IdTypeMap<type_id::TIMESTAMP_DAYS>::type>(
        std::forward<Ts>(args)...);
    case type_id::TIMESTAMP_SECONDS:
      return f.template operator()<typename IdTypeMap<type_id::TIMESTAMP_SECONDS>::type>(
        std::forward<Ts>(args)...);
    case type_id::TIMESTAMP_MILLISECONDS:
      return f.template operator()<typename IdTypeMap<type_id::TIMESTAMP_MILLISECONDS>::type>(
        std::forward<Ts>(args)...);
    case type_id::TIMESTAMP_MICROSECONDS:
      return f.template operator()<typename IdTypeMap<type_id::TIMESTAMP_MICROSECONDS>::type>(
        std::forward<Ts>(args)...);
    case type_id::TIMESTAMP_NANOSECONDS:
      return f.template operator()<typename IdTypeMap<type_id::TIMESTAMP_NANOSECONDS>::type>(
        std::forward<Ts>(args)...);
    case type_id::DURATION_DAYS:
      return f.template operator()<typename IdTypeMap<type_id::DURATION_DAYS>::type>(
        std::forward<Ts>(args)...);
    case type_id::DURATION_SECONDS:
      return f.template operator()<typename IdTypeMap<type_id::DURATION_SECONDS>::type>(
        std::forward<Ts>(args)...);
    case type_id::DURATION_MILLISECONDS:
      return f.template operator()<typename IdTypeMap<type_id::DURATION_MILLISECONDS>::type>(
        std::forward<Ts>(args)...);
    case type_id::DURATION_MICROSECONDS:
      return f.template operator()<typename IdTypeMap<type_id::DURATION_MICROSECONDS>::type>(
        std::forward<Ts>(args)...);
    case type_id::DURATION_NANOSECONDS:
      return f.template operator()<typename IdTypeMap<type_id::DURATION_NANOSECONDS>::type>(
        std::forward<Ts>(args)...);
    case type_id::DICTIONARY32:
      return f.template operator()<typename IdTypeMap<type_id::DICTIONARY32>::type>(
        std::forward<Ts>(args)...);
    case type_id::LIST:
      return f.template operator()<typename IdTypeMap<type_id::LIST>::type>(
        std::forward<Ts>(args)...);
    case type_id::DECIMAL32:
      return f.template operator()<typename IdTypeMap<type_id::DECIMAL32>::type>(
        std::forward<Ts>(args)...);
    case type_id::DECIMAL64:
      return f.template operator()<typename IdTypeMap<type_id::DECIMAL64>::type>(
        std::forward<Ts>(args)...);
    case type_id::DECIMAL128:
      return f.template operator()<typename IdTypeMap<type_id::DECIMAL128>::type>(
        std::forward<Ts>(args)...);
    case type_id::STRUCT:
      return f.template operator()<typename IdTypeMap<type_id::STRUCT>::type>(
        std::forward<Ts>(args)...);
    default: {
#ifndef __CUDA_ARCH__
      CUDF_FAIL("Invalid type_id.");
#else
      CUDF_UNREACHABLE("Invalid type_id.");
#endif
    }
  }
}
 
namespace detail {
template <typename T1>
struct double_type_dispatcher_second_type {
#pragma nv_exec_check_disable
  template <typename T2, typename F, typename... Ts>
  CUDF_HOST_DEVICE __forceinline__ decltype(auto) operator()(F&& f, Ts&&... args) const
  {
    return f.template operator()<T1, T2>(std::forward<Ts>(args)...);
  }
};
 
template <template <cudf::type_id> typename IdTypeMap>
struct double_type_dispatcher_first_type {
#pragma nv_exec_check_disable
  template <typename T1, typename F, typename... Ts>
  CUDF_HOST_DEVICE __forceinline__ decltype(auto) operator()(cudf::data_type type2,
                                                             F&& f,
                                                             Ts&&... args) const
  {
    return type_dispatcher<IdTypeMap>(type2,
                                      detail::double_type_dispatcher_second_type<T1>{},
                                      std::forward<F>(f),
                                      std::forward<Ts>(args)...);
  }
};
}  // namespace detail
 
#pragma nv_exec_check_disable
template <template <cudf::type_id> typename IdTypeMap = id_to_type_impl, typename F, typename... Ts>
CUDF_HOST_DEVICE __forceinline__ constexpr decltype(auto) double_type_dispatcher(
  cudf::data_type type1, cudf::data_type type2, F&& f, Ts&&... args)
{
  return type_dispatcher<IdTypeMap>(type1,
                                    detail::double_type_dispatcher_first_type<IdTypeMap>{},
                                    type2,
                                    std::forward<F>(f),
                                    std::forward<Ts>(args)...);
}
  // end of group
}  // namespace cudf