diff --git a/CHANGELOG.md b/CHANGELOG.md index 1176df005fd..44e52c98b23 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -50,6 +50,7 @@ - PR #5658 Add `filter_tokens` nvtext API - PR #5666 Add `filter_characters_of_type` strings API - PR #5673 Always build and test with per-thread default stream enabled in the GPU CI build +- PR #5704 Initial `fixed_point` Column Support - PR #5716 Add `double_type_dispatcher` to libcudf - PR #5739 Add `nvtext::detokenize` API - PR #5645 Enforce pd.NA and Pandas nullable dtype parity diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 2f84663bcf3..94cf9486e96 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -623,6 +623,7 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include ${CMAKE_BINARY_DIR}/include/jit/types.hpp.jit ${CMAKE_BINARY_DIR}/include/bit.hpp.jit ${CMAKE_BINARY_DIR}/include/jit/timestamps.hpp.jit + ${CMAKE_BINARY_DIR}/include/jit/fixed_point.hpp.jit ${CMAKE_BINARY_DIR}/include/jit/durations.hpp.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/details/__config.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/limits.jit @@ -632,6 +633,8 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/ratio.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/type_traits.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/version.jit + ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cmath.jit + ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cassert.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__config.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__undef_macros.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cfloat.jit @@ -640,10 +643,13 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/limits.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/ratio.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/type_traits.jit + ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cmath.jit + ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cassert.jit MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/types.h ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/types.hpp ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/utilities/bit.hpp ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/wrappers/timestamps.hpp + ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/fixed_point/fixed_point.hpp ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/wrappers/durations.hpp ${LIBCUDACXX_INCLUDE_DIR}/details/__config ${LIBCUDACXX_INCLUDE_DIR}/simt/limits @@ -653,6 +659,8 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include ${LIBCUDACXX_INCLUDE_DIR}/simt/ratio ${LIBCUDACXX_INCLUDE_DIR}/simt/type_traits ${LIBCUDACXX_INCLUDE_DIR}/simt/version + ${LIBCUDACXX_INCLUDE_DIR}/simt/cmath + ${LIBCUDACXX_INCLUDE_DIR}/simt/cassert ${LIBCXX_INCLUDE_DIR}/__config ${LIBCXX_INCLUDE_DIR}/__undef_macros ${LIBCXX_INCLUDE_DIR}/cfloat @@ -661,6 +669,8 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include ${LIBCXX_INCLUDE_DIR}/limits ${LIBCXX_INCLUDE_DIR}/ratio ${LIBCXX_INCLUDE_DIR}/type_traits + ${LIBCXX_INCLUDE_DIR}/cmath + ${LIBCXX_INCLUDE_DIR}/cassert # stringified headers are placed underneath the bin include jit directory and end in ".jit" COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/types.h > ${CMAKE_BINARY_DIR}/include/jit/types.h.jit @@ -668,6 +678,7 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/utilities/bit.hpp > ${CMAKE_BINARY_DIR}/include/bit.hpp.jit COMMAND ${CMAKE_BINARY_DIR}/stringify ../src/rolling/rolling_jit_detail.hpp > ${CMAKE_BINARY_DIR}/include/rolling_jit_detail.hpp.jit COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/wrappers/timestamps.hpp > ${CMAKE_BINARY_DIR}/include/jit/timestamps.hpp.jit + COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/fixed_point/fixed_point.hpp > ${CMAKE_BINARY_DIR}/include/jit/fixed_point.hpp.jit COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/wrappers/durations.hpp > ${CMAKE_BINARY_DIR}/include/jit/durations.hpp.jit COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/details/__config libcudacxx_details_config > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/details/__config.jit COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/limits libcudacxx_simt_limits > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/limits.jit @@ -676,6 +687,8 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/ctime libcudacxx_simt_ctime > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/ctime.jit COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/ratio libcudacxx_simt_ratio > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/ratio.jit COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/type_traits libcudacxx_simt_type_traits > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/type_traits.jit + COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/cmath libcudacxx_simt_cmath > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cmath.jit + COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/cassert libcudacxx_simt_cassert > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cassert.jit COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/version libcudacxx_simt_version > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/version.jit COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/__config libcxx_config > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__config.jit COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/__undef_macros libcxx_undef_macros > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__undef_macros.jit @@ -685,6 +698,8 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/limits libcxx_limits > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/limits.jit COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/ratio libcxx_ratio > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/ratio.jit COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/type_traits libcxx_type_traits > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/type_traits.jit + COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/cmath libcxx_cmath > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cmath.jit + COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/cassert libcxx_cassert > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cassert.jit ) add_custom_target(stringify_run DEPENDS @@ -692,6 +707,7 @@ add_custom_target(stringify_run DEPENDS ${CMAKE_BINARY_DIR}/include/jit/types.hpp.jit ${CMAKE_BINARY_DIR}/include/bit.hpp.jit ${CMAKE_BINARY_DIR}/include/jit/timestamps.hpp.jit + ${CMAKE_BINARY_DIR}/include/jit/fixed_point.hpp.jit ${CMAKE_BINARY_DIR}/include/jit/durations.hpp.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/details/__config.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/limits.jit @@ -700,6 +716,8 @@ add_custom_target(stringify_run DEPENDS ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/ctime.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/ratio.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/type_traits.jit + ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cmath.jit + ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cassert.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/version.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__config.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__undef_macros.jit @@ -708,7 +726,10 @@ add_custom_target(stringify_run DEPENDS ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/ctime.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/limits.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/ratio.jit - ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/type_traits.jit) + ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/type_traits.jit + ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cmath.jit + ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cassert.jit + ) add_dependencies(cudf stringify_run) diff --git a/cpp/include/cudf/column/column_factories.hpp b/cpp/include/cudf/column/column_factories.hpp index 8738f7ce56b..3fc4c60b51c 100644 --- a/cpp/include/cudf/column/column_factories.hpp +++ b/cpp/include/cudf/column/column_factories.hpp @@ -86,7 +86,7 @@ std::unique_ptr make_numeric_column( cudaStream_t stream = 0, rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource()) { - CUDF_EXPECTS(is_numeric(type), "Invalid, non-numeric type."); + CUDF_EXPECTS(is_numeric(type) || is_fixed_point(type), "Invalid, non-numeric type."); return std::make_unique(type, size, rmm::device_buffer{size * cudf::size_of(type), stream, mr}, diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index 0be3400faf8..48716cf447d 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.cuh +++ b/cpp/include/cudf/detail/aggregation/aggregation.cuh @@ -96,11 +96,12 @@ struct update_target_element { }; template -struct update_target_element()>> { +struct update_target_element< + Source, + aggregation::MIN, + target_has_nulls, + source_has_nulls, + std::enable_if_t() && !is_fixed_point()>> { __device__ void operator()(mutable_column_device_view target, size_type target_index, column_device_view source, @@ -117,11 +118,12 @@ struct update_target_element -struct update_target_element()>> { +struct update_target_element< + Source, + aggregation::MAX, + target_has_nulls, + source_has_nulls, + std::enable_if_t() && !is_fixed_point()>> { __device__ void operator()(mutable_column_device_view target, size_type target_index, column_device_view source, @@ -138,11 +140,12 @@ struct update_target_element -struct update_target_element()>> { +struct update_target_element< + Source, + aggregation::SUM, + target_has_nulls, + source_has_nulls, + std::enable_if_t() && !is_fixed_point()>> { __device__ void operator()(mutable_column_device_view target, size_type target_index, column_device_view source, @@ -359,7 +362,7 @@ struct identity_initializer { template static constexpr bool is_supported() { - return cudf::is_fixed_width() and + return cudf::is_fixed_width() && !is_fixed_point() and (k == aggregation::SUM or k == aggregation::MIN or k == aggregation::MAX or k == aggregation::COUNT_VALID or k == aggregation::COUNT_ALL or k == aggregation::ARGMAX or k == aggregation::ARGMIN); diff --git a/cpp/include/cudf/detail/reduction.cuh b/cpp/include/cudf/detail/reduction.cuh index 15a852c4539..84cde38fab8 100644 --- a/cpp/include/cudf/detail/reduction.cuh +++ b/cpp/include/cudf/detail/reduction.cuh @@ -79,9 +79,8 @@ std::unique_ptr reduce(InputIterator d_in, identity, stream); - using ScalarType = cudf::scalar_type_t; - auto s = new ScalarType( - std::move(dev_result), true, stream, mr); // only for string_view, data is copied + // only for string_view, data is copied + auto s = new cudf::scalar_type_t(std::move(dev_result), true, stream, mr); return std::unique_ptr(s); } diff --git a/cpp/include/cudf/detail/utilities/device_operators.cuh b/cpp/include/cudf/detail/utilities/device_operators.cuh index 4de52d3ca1d..460b0181abe 100644 --- a/cpp/include/cudf/detail/utilities/device_operators.cuh +++ b/cpp/include/cudf/detail/utilities/device_operators.cuh @@ -27,6 +27,10 @@ #include #include #include + +// will fail to compile if grouped with the includes above +#include + #include namespace cudf { @@ -40,11 +44,18 @@ struct DeviceSum { return lhs + rhs; } - template + template ()>* = nullptr> static constexpr T identity() { return T{0}; } + + template ()>* = nullptr> + static constexpr T identity() + { + CUDF_FAIL("fixed_point does not yet support device operator identity"); + return T{}; + } }; /* @brief `count` operator - used in rolling windows */ @@ -88,9 +99,17 @@ struct DeviceMin { } template ::value>* = nullptr> + typename std::enable_if_t::value && + !cudf::is_fixed_point()>* = nullptr> + static constexpr T identity() + { + return std::numeric_limits::max(); + } + + template ()>* = nullptr> static constexpr T identity() { + CUDF_FAIL("fixed_point does not yet support DeviceMin identity"); return std::numeric_limits::max(); } @@ -118,11 +137,20 @@ struct DeviceMax { } template ::value>* = nullptr> + typename std::enable_if_t::value && + !cudf::is_fixed_point()>* = nullptr> + static constexpr T identity() + { + return std::numeric_limits::lowest(); + } + + template ()>* = nullptr> static constexpr T identity() { + CUDF_FAIL("fixed_point does not yet support DeviceMax identity"); return std::numeric_limits::lowest(); } + template ::value>* = nullptr> CUDA_HOST_DEVICE_CALLABLE static constexpr T identity() @@ -145,11 +173,18 @@ struct DeviceProduct { return lhs * rhs; } - template + template ()>* = nullptr> static constexpr T identity() { return T{1}; } + + template ()>* = nullptr> + static constexpr T identity() + { + CUDF_FAIL("fixed_point does not yet support DeviceProduct identity"); + return T{1, numeric::scale_type{0}}; + } }; /* @brief binary `and` operator */ diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index 7063cdeec2d..3634d86f37a 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -14,11 +14,20 @@ * limitations under the License. */ +#pragma once + +#define _LIBCUDACXX_USE_CXX17_TYPE_TRAITS + +// Note: The versions are used in order for Jitify to work with our fixed_point type. +// Jitify is needed for several algorithms (binaryop, rolling, etc) +#include +#include // add simt namespace + +#include #include #include -#include -#include -#include + +#include //! `fixed_point` and supporting types namespace numeric { @@ -51,13 +60,13 @@ enum class Radix : int32_t { BASE_2 = 2, BASE_10 = 10 }; template constexpr inline auto is_supported_representation_type() { - return std::is_same::value || std::is_same::value; + return simt::std::is_same::value || simt::std::is_same::value; } template constexpr inline auto is_supported_construction_value_type() { - return std::is_integral::value || std::is_floating_point::value; + return simt::std::is_integral::value || simt::std::is_floating_point::value; } // Helper functions for `fixed_point` type @@ -78,8 +87,8 @@ namespace detail { template ::value && - is_supported_representation_type())>* = nullptr> + typename simt::std::enable_if_t<(simt::std::is_same::value && + is_supported_representation_type())>* = nullptr> CUDA_HOST_DEVICE_CALLABLE Rep ipow(T exponent) { if (exponent == 0) return static_cast(1); @@ -200,7 +209,7 @@ CUDA_HOST_DEVICE_CALLABLE constexpr T shift(T const& val, scale_type const& scal template ::value>* = nullptr> + typename simt::std::enable_if_t::value>* = nullptr> CUDA_HOST_DEVICE_CALLABLE auto shift_with_precise_round(T const& value, scale_type const& scale) -> Rep { @@ -228,7 +237,7 @@ CUDA_HOST_DEVICE_CALLABLE auto shift_with_precise_round(T const& value, scale_ty template ::value>* = nullptr> + typename simt::std::enable_if_t::value>* = nullptr> CUDA_HOST_DEVICE_CALLABLE auto shift_with_precise_round(T const& value, scale_type const& scale) -> Rep { @@ -256,7 +265,7 @@ CUDA_HOST_DEVICE_CALLABLE auto shift_with_precise_round(T const& value, scale_ty * @tparam Rep The representation type (either `int32_t` or `int64_t`) */ template ()>* = nullptr> + typename simt::std::enable_if_t()>* = nullptr> struct scaled_integer { Rep value; scale_type scale; @@ -286,8 +295,8 @@ class fixed_point { * @param scale The exponent that is applied to Rad to perform shifting */ template () && - is_supported_representation_type()>* = nullptr> + typename simt::std::enable_if_t() && + is_supported_representation_type()>* = nullptr> CUDA_HOST_DEVICE_CALLABLE explicit fixed_point(T const& value, scale_type const& scale) : _value{detail::shift_with_precise_round(value, scale)}, _scale{scale} { @@ -316,7 +325,7 @@ class fixed_point { * @return The `fixed_point` number in base 10 (aka human readable format) */ template ()>* = nullptr> + typename simt::std::enable_if_t()>* = nullptr> CUDA_HOST_DEVICE_CALLABLE explicit constexpr operator U() const { return detail::shift(static_cast(_value), detail::negate(_scale)); @@ -562,9 +571,9 @@ class fixed_point { template std::string print_rep() { - if (std::is_same::value) + if (simt::std::is_same::value) return "int32_t"; - else if (std::is_same::value) + else if (simt::std::is_same::value) return "int64_t"; else return "unknown type"; @@ -581,8 +590,8 @@ std::string print_rep() template CUDA_HOST_DEVICE_CALLABLE auto addition_overflow(T lhs, T rhs) { - return rhs > 0 ? lhs > std::numeric_limits::max() - rhs - : lhs < std::numeric_limits::min() - rhs; + return rhs > 0 ? lhs > simt::std::numeric_limits::max() - rhs + : lhs < simt::std::numeric_limits::min() - rhs; } /** @brief Function for identifying integer overflow when subtracting @@ -596,8 +605,8 @@ CUDA_HOST_DEVICE_CALLABLE auto addition_overflow(T lhs, T rhs) template CUDA_HOST_DEVICE_CALLABLE auto subtraction_overflow(T lhs, T rhs) { - return rhs > 0 ? lhs < std::numeric_limits::min() + rhs - : lhs > std::numeric_limits::max() + rhs; + return rhs > 0 ? lhs < simt::std::numeric_limits::min() + rhs + : lhs > simt::std::numeric_limits::max() + rhs; } /** @brief Function for identifying integer overflow when dividing @@ -611,7 +620,7 @@ CUDA_HOST_DEVICE_CALLABLE auto subtraction_overflow(T lhs, T rhs) template CUDA_HOST_DEVICE_CALLABLE auto division_overflow(T lhs, T rhs) { - return lhs == std::numeric_limits::min() && rhs == -1; + return lhs == simt::std::numeric_limits::min() && rhs == -1; } /** @brief Function for identifying integer overflow when multiplying @@ -625,8 +634,8 @@ CUDA_HOST_DEVICE_CALLABLE auto division_overflow(T lhs, T rhs) template CUDA_HOST_DEVICE_CALLABLE auto multiplication_overflow(T lhs, T rhs) { - auto const min = std::numeric_limits::min(); - auto const max = std::numeric_limits::max(); + auto const min = simt::std::numeric_limits::min(); + auto const max = simt::std::numeric_limits::max(); if (rhs > 0) return lhs > max / rhs || lhs < min / rhs; else if (rhs < -1) @@ -757,19 +766,8 @@ CUDA_HOST_DEVICE_CALLABLE bool operator>(fixed_point const& lhs, return lhs.rescaled(scale)._value > rhs.rescaled(scale)._value; } -/** - * @brief ostream operator for outputting `fixed_point` number - * - * @tparam Rep Representation type of number being output - * @tparam Rad Radix (base) type of number being output - * @param os target ostream - * @return fp `fixed_point` number being output - */ -template -std::ostream& operator<<(std::ostream& os, fixed_point const& fp) -{ - return os << static_cast(fp); -} +using decimal32 = fixed_point; +using decimal64 = fixed_point; /** @} */ // end of group } // namespace numeric diff --git a/cpp/include/cudf/scalar/scalar.hpp b/cpp/include/cudf/scalar/scalar.hpp index f6becdfd1ae..d647297dc02 100644 --- a/cpp/include/cudf/scalar/scalar.hpp +++ b/cpp/include/cudf/scalar/scalar.hpp @@ -24,7 +24,6 @@ #include #include -#include #include #include @@ -94,10 +93,10 @@ class scalar { * @note Do not use this constructor directly. Instead, use a factory method * like make_numeric_scalar or make_string_scalar * - * @param type Data type of the scalar - * @param is_valid Whether the value held by the scalar is valid - * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource to use for device memory allocation + * @param[in] type Data type of the scalar + * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation */ scalar(data_type type, bool is_valid = false, @@ -164,10 +163,10 @@ class fixed_width_scalar : public scalar { /** * @brief Construct a new fixed width scalar object * - * @param value The initial value of the scalar - * @param is_valid Whether the value held by the scalar is valid - * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource to use for device memory allocation + * @param[in] value The initial value of the scalar + * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation */ fixed_width_scalar(T value, bool is_valid = true, @@ -182,6 +181,8 @@ class fixed_width_scalar : public scalar { * * @param[in] data The scalar's data in device memory * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation */ fixed_width_scalar(rmm::device_scalar&& data, bool is_valid = true, @@ -217,10 +218,10 @@ class numeric_scalar : public detail::fixed_width_scalar { /** * @brief Construct a new numeric scalar object * - * @param value The initial value of the scalar - * @param is_valid Whether the value held by the scalar is valid - * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource to use for device memory allocation + * @param[in] value The initial value of the scalar + * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation */ numeric_scalar(T value, bool is_valid = true, @@ -235,6 +236,8 @@ class numeric_scalar : public detail::fixed_width_scalar { * * @param[in] data The scalar's data in device memory * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation */ numeric_scalar(rmm::device_scalar&& data, bool is_valid = true, @@ -245,6 +248,58 @@ class numeric_scalar : public detail::fixed_width_scalar { } }; +/** + * @brief An owning class to represent a fixed_point number in device memory + * + * @ingroup scalar_classes + * + * @tparam T the data type of the fixed_point number + */ +template +class fixed_point_scalar : public detail::fixed_width_scalar { + static_assert(is_fixed_point(), "Unexpected non-fixed_point type."); + + public: + fixed_point_scalar() = default; + ~fixed_point_scalar() = default; + fixed_point_scalar(fixed_point_scalar&& other) = default; + fixed_point_scalar(fixed_point_scalar const& other) = default; + fixed_point_scalar& operator=(fixed_point_scalar const& other) = delete; + fixed_point_scalar& operator=(fixed_point_scalar&& other) = delete; + + /** + * @brief Construct a new fixed_point scalar object + * + * @param[in] value The initial value of the scalar + * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation + */ + fixed_point_scalar(T value, + bool is_valid = true, + cudaStream_t stream = 0, + rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource()) + : detail::fixed_width_scalar(value, is_valid, stream, mr) + { + } + + /** + * @brief Construct a new fixed_point scalar object from existing device memory. + * + * @param[in] data The scalar's data in device memory + * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation + */ + fixed_point_scalar(rmm::device_scalar&& data, + bool is_valid = true, + cudaStream_t stream = 0, + rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource()) + : detail::fixed_width_scalar(std::forward>(data), is_valid, stream, mr) + { + } +}; + /** * @brief An owning class to represent a string in device memory * @@ -264,10 +319,10 @@ class string_scalar : public scalar { /** * @brief Construct a new string scalar object * - * @param value The value of the string - * @param is_valid Whether the value held by the scalar is valid - * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource to use for device memory allocation + * @param[in] value The value of the string + * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation */ string_scalar(std::string const& string, bool is_valid = true, @@ -281,10 +336,10 @@ class string_scalar : public scalar { * @brief Construct a new string scalar object from string_view * Note that this function copies the data pointed by string_view. * - * @param source string_view pointing string value to copy - * @param is_valid Whether the value held by the scalar is valid - * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource to use for device memory allocation + * @param[in] source string_view pointing string value to copy + * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation */ string_scalar(value_type const& source, bool is_valid = true, @@ -299,10 +354,10 @@ class string_scalar : public scalar { * @brief Construct a new string scalar object from string_view in device memory * Note that this function copies the data pointed by string_view. * - * @param data device_scalar string_view pointing string value to copy - * @param is_valid Whether the value held by the scalar is valid - * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource to use for device memory allocation + * @param[in] data device_scalar string_view pointing string value to copy + * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation */ string_scalar(rmm::device_scalar& data, bool is_valid = true, @@ -368,10 +423,10 @@ class chrono_scalar : public detail::fixed_width_scalar { /** * @brief Construct a new chrono scalar object * - * @param value The initial value of the scalar - * @param is_valid Whether the value held by the scalar is valid - * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource to use for device memory allocation + * @param[in] value The initial value of the scalar + * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation */ chrono_scalar(T value, bool is_valid = true, @@ -384,10 +439,10 @@ class chrono_scalar : public detail::fixed_width_scalar { /** * @brief Construct a new chrono scalar object from an integer * - * @param value Integer representing number of ticks since the UNIX epoch - * @param is_valid Whether the value held by the scalar is valid - * @param stream CUDA stream used for device memory operations. - * @param mr Device memory resource to use for device memory allocation + * @param[in] value Integer representing number of ticks since the UNIX epoch + * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation */ chrono_scalar(typename T::rep value, bool is_valid, @@ -402,6 +457,8 @@ class chrono_scalar : public detail::fixed_width_scalar { * * @param[in] data The scalar's data in device memory * @param[in] is_valid Whether the value held by the scalar is valid + * @param[in] stream CUDA stream used for device memory operations. + * @param[in] mr Device memory resource to use for device memory allocation */ chrono_scalar(rmm::device_scalar&& data, bool is_valid = true, diff --git a/cpp/include/cudf/scalar/scalar_device_view.cuh b/cpp/include/cudf/scalar/scalar_device_view.cuh index ddc77c6003a..3c80daaa07b 100644 --- a/cpp/include/cudf/scalar/scalar_device_view.cuh +++ b/cpp/include/cudf/scalar/scalar_device_view.cuh @@ -122,6 +122,18 @@ class numeric_scalar_device_view : public detail::fixed_width_scalar_device_view } }; +/** + * @brief A type of scalar_device_view that stores a pointer to a fixed_point value + */ +template +class fixed_point_scalar_device_view : public detail::fixed_width_scalar_device_view { + public: + fixed_point_scalar_device_view(data_type type, T* data, bool* is_valid) + : detail::fixed_width_scalar_device_view(type, data, is_valid) + { + } +}; + /** * @brief A type of scalar_device_view that stores a pointer to a string value */ diff --git a/cpp/include/cudf/types.hpp b/cpp/include/cudf/types.hpp index 5bce5dc393c..f19ce096efb 100644 --- a/cpp/include/cudf/types.hpp +++ b/cpp/include/cudf/types.hpp @@ -63,6 +63,10 @@ class list_view; class scalar; template class numeric_scalar; + +template +class fixed_point_scalar; + class string_scalar; template class timestamp_scalar; @@ -71,6 +75,10 @@ class duration_scalar; template class numeric_scalar_device_view; + +template +class fixed_point_scalar_device_view; + class string_scalar_device_view; template class timestamp_scalar_device_view; @@ -207,6 +215,8 @@ enum class type_id : int32_t { DICTIONARY32, ///< Dictionary type using int32 indices STRING, ///< String elements LIST, ///< List elements + DECIMAL32, ///< Fixed-point type with int32_t + DECIMAL64, ///< Fixed-point type with int64_t // `NUM_TYPE_IDS` must be last! NUM_TYPE_IDS ///< Total number of type ids }; diff --git a/cpp/include/cudf/utilities/traits.hpp b/cpp/include/cudf/utilities/traits.hpp index 637483ef747..d99d5daf10a 100644 --- a/cpp/include/cudf/utilities/traits.hpp +++ b/cpp/include/cudf/utilities/traits.hpp @@ -23,7 +23,7 @@ #include #include -#include +#include namespace cudf { @@ -329,6 +329,39 @@ constexpr inline bool is_timestamp(data_type type) return cudf::type_dispatcher(type, is_timestamp_impl{}); } +/** + * @brief Indicates whether the type `T` is a fixed-point type. + * + * @tparam T The type to verify + * @return true `T` is a fixed-point type + * @return false `T` is not a fixed-point type + **/ +template +constexpr inline bool is_fixed_point() +{ + return std::is_same::value || std::is_same::value; +} + +struct is_fixed_point_impl { + template + bool operator()() + { + return is_fixed_point(); + } +}; + +/** + * @brief Indicates whether `type` is a fixed point `data_type`. + * + * @param type The `data_type` to verify + * @return true `type` is a fixed point type + * @return false `type` is not a fixed point type + **/ +constexpr inline bool is_fixed_point(data_type type) +{ + return cudf::type_dispatcher(type, is_fixed_point_impl{}); +} + /** * @brief Indicates whether the type `T` is a duration type. * @@ -414,7 +447,7 @@ constexpr inline bool is_fixed_width() { // TODO Add fixed width wrapper types // Is a category fixed width? - return cudf::is_numeric() || cudf::is_chrono(); + return cudf::is_numeric() || cudf::is_chrono() || cudf::is_fixed_point(); } struct is_fixed_width_impl { diff --git a/cpp/include/cudf/utilities/type_dispatcher.hpp b/cpp/include/cudf/utilities/type_dispatcher.hpp index 4a18f6bb7b1..47041084ba9 100644 --- a/cpp/include/cudf/utilities/type_dispatcher.hpp +++ b/cpp/include/cudf/utilities/type_dispatcher.hpp @@ -24,6 +24,8 @@ #include #include +#include + /** * @file type_dispatcher.hpp * @brief Defines the mapping between `cudf::type_id` runtime type information @@ -132,6 +134,8 @@ 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); template struct type_to_scalar_type_impl { @@ -171,6 +175,18 @@ struct type_to_scalar_type_impl { using ScalarDeviceType = cudf::string_scalar_device_view; }; +template <> +struct type_to_scalar_type_impl { + using ScalarType = cudf::fixed_point_scalar; + using ScalarDeviceType = cudf::fixed_point_scalar_device_view; +}; + +template <> +struct type_to_scalar_type_impl { + using ScalarType = cudf::fixed_point_scalar; + using ScalarDeviceType = cudf::fixed_point_scalar_device_view; +}; + template <> // TODO: this is a temporary solution for make_pair_iterator struct type_to_scalar_type_impl { using ScalarType = cudf::numeric_scalar; @@ -400,6 +416,12 @@ CUDA_HOST_DEVICE_CALLABLE constexpr decltype(auto) type_dispatcher(cudf::data_ty case type_id::LIST: return f.template operator()::type>( std::forward(args)...); + case type_id::DECIMAL32: + return f.template operator()::type>( + std::forward(args)...); + case type_id::DECIMAL64: + return f.template operator()::type>( + std::forward(args)...); default: { #ifndef __CUDA_ARCH__ CUDF_FAIL("Unsupported type_id."); diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 66c3389723d..0cdd2db67cb 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -37,6 +37,7 @@ #include #include +#include #include #include #include @@ -75,7 +76,8 @@ const std::vector header_names{"operation.h", cudf_types_hpp, cudf_utilities_bit_hpp, cudf_wrappers_timestamps_hpp, - cudf_wrappers_durations_hpp}; + cudf_wrappers_durations_hpp, + cudf_fixed_point_fixed_point_hpp}; std::istream* headers_code(std::string filename, std::iostream& stream) { @@ -330,14 +332,15 @@ std::unique_ptr binary_operation(column_view const& lhs, CUDF_EXPECTS(is_fixed_width(lhs.type()), "Invalid/Unsupported lhs datatype"); CUDF_EXPECTS(is_fixed_width(rhs.type()), "Invalid/Unsupported rhs datatype"); - std::unique_ptr out; - if (binops::null_using_binop(op)) { - out = make_fixed_width_column(output_type, rhs.size(), mask_state::ALL_VALID, stream, mr); - } else { - auto new_mask = bitmask_and(table_view({lhs, rhs}), mr, stream); - out = make_fixed_width_column( - output_type, lhs.size(), std::move(new_mask), cudf::UNKNOWN_NULL_COUNT, stream, mr); - } + std::unique_ptr out = [&] { + if (binops::null_using_binop(op)) { + return make_fixed_width_column(output_type, rhs.size(), mask_state::ALL_VALID, stream, mr); + } else { + auto new_mask = bitmask_and(table_view({lhs, rhs}), mr, stream); + return make_fixed_width_column( + output_type, lhs.size(), std::move(new_mask), cudf::UNKNOWN_NULL_COUNT, stream, mr); + } + }(); // Check for 0 sized data if (lhs.size() == 0 || rhs.size() == 0) { return out; } diff --git a/cpp/src/binaryop/jit/code/kernel.cpp b/cpp/src/binaryop/jit/code/kernel.cpp index bb7cec6e404..ba7ab305de4 100644 --- a/cpp/src/binaryop/jit/code/kernel.cpp +++ b/cpp/src/binaryop/jit/code/kernel.cpp @@ -31,6 +31,7 @@ const char* kernel = #include #include #include + #include #include "operation.h" template diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu index 1eb41ffae15..2def3471729 100644 --- a/cpp/src/copying/copy.cu +++ b/cpp/src/copying/copy.cu @@ -20,6 +20,7 @@ #include #include #include +#include "cudf/fixed_point/fixed_point.hpp" namespace cudf { namespace detail { @@ -115,6 +116,42 @@ struct copy_if_else_functor_impl { } }; +/** + * @brief Specialization of copy_if_else_functor for decimal32. + */ +template +struct copy_if_else_functor_impl { + std::unique_ptr operator()(Left const& lhs, + Right const& rhs, + size_type size, + bool left_nullable, + bool right_nullable, + Filter filter, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) + { + CUDF_FAIL("copy_if_else not supported for decimal32 yet"); + } +}; + +/** + * @brief Specialization of copy_if_else_functor for decimal64. + */ +template +struct copy_if_else_functor_impl { + std::unique_ptr operator()(Left const& lhs, + Right const& rhs, + size_type size, + bool left_nullable, + bool right_nullable, + Filter filter, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) + { + CUDF_FAIL("copy_if_else not supported for decimal64 yet"); + } +}; + /** * @brief Functor called by the `type_dispatcher` to invoke copy_if_else on combinations * of column_view and scalar diff --git a/cpp/src/copying/get_element.cu b/cpp/src/copying/get_element.cu index 04edae9366b..6408d1c7720 100644 --- a/cpp/src/copying/get_element.cu +++ b/cpp/src/copying/get_element.cu @@ -28,7 +28,7 @@ namespace detail { namespace { struct get_element_functor { - template ()> *p = nullptr> + template () && !is_fixed_point()> *p = nullptr> std::unique_ptr operator()( column_view const &input, size_type index, @@ -114,6 +114,26 @@ struct get_element_functor { { CUDF_FAIL("get_element_functor not supported for list_view"); } + + template ::value> *p = nullptr> + std::unique_ptr operator()( + column_view const &input, + size_type index, + cudaStream_t stream = 0, + rmm::mr::device_memory_resource *mr = rmm::mr::get_default_resource()) + { + CUDF_FAIL("get_element_functor not supported for decimal32"); + } + + template ::value> *p = nullptr> + std::unique_ptr operator()( + column_view const &input, + size_type index, + cudaStream_t stream = 0, + rmm::mr::device_memory_resource *mr = rmm::mr::get_default_resource()) + { + CUDF_FAIL("get_element_functor not supported for decimal64"); + } }; } // namespace diff --git a/cpp/src/io/csv/csv_gpu.cu b/cpp/src/io/csv/csv_gpu.cu index 3eb53f422d0..9d91ebb95f0 100644 --- a/cpp/src/io/csv/csv_gpu.cu +++ b/cpp/src/io/csv/csv_gpu.cu @@ -20,6 +20,7 @@ #include "datetime.cuh" #include +#include #include #include #include @@ -402,6 +403,28 @@ __inline__ __device__ cudf::list_view decode_value(const char *data, return cudf::list_view{}; } +// The purpose of this is merely to allow compilation ONLY +// TODO : make this work for csv +template <> +__inline__ __device__ numeric::decimal32 decode_value(const char *data, + long start, + long end, + ParseOptions const &opts) +{ + return numeric::decimal32{}; +} + +// The purpose of this is merely to allow compilation ONLY +// TODO : make this work for csv +template <> +__inline__ __device__ numeric::decimal64 decode_value(const char *data, + long start, + long end, + ParseOptions const &opts) +{ + return numeric::decimal64{}; +} + /** * @brief Functor for converting CSV raw data to typed value. */ diff --git a/cpp/src/io/json/json_gpu.cu b/cpp/src/io/json/json_gpu.cu index d15b4465f47..1347f2ecef5 100644 --- a/cpp/src/io/json/json_gpu.cu +++ b/cpp/src/io/json/json_gpu.cu @@ -27,6 +27,7 @@ #include #include +#include #include #include @@ -252,6 +253,7 @@ __inline__ __device__ cudf::string_view decode_value(const char *data, { return cudf::string_view{}; } + template <> __inline__ __device__ cudf::dictionary32 decode_value(const char *data, uint64_t start, @@ -260,6 +262,7 @@ __inline__ __device__ cudf::dictionary32 decode_value(const char *data, { return cudf::dictionary32{}; } + template <> __inline__ __device__ cudf::list_view decode_value(const char *data, uint64_t start, @@ -269,6 +272,24 @@ __inline__ __device__ cudf::list_view decode_value(const char *data, return cudf::list_view{}; } +template <> +__inline__ __device__ numeric::decimal32 decode_value(const char *data, + uint64_t start, + uint64_t end, + ParseOptions const &opts) +{ + return numeric::decimal32{}; +} + +template <> +__inline__ __device__ numeric::decimal64 decode_value(const char *data, + uint64_t start, + uint64_t end, + ParseOptions const &opts) +{ + return numeric::decimal64{}; +} + /** * @brief Functor for converting plain text data to cuDF data type value. */ diff --git a/cpp/src/io/utilities/parsing_utils.cuh b/cpp/src/io/utilities/parsing_utils.cuh index 5867c0d277a..5e48365936a 100644 --- a/cpp/src/io/utilities/parsing_utils.cuh +++ b/cpp/src/io/utilities/parsing_utils.cuh @@ -143,7 +143,7 @@ template __inline__ __device__ T parse_numeric(const char* data, long start, long end, ParseOptions const& opts) { - T value = 0; + T value{}; bool all_digits_valid = true; // Handle negative values if necessary diff --git a/cpp/src/jit/common_headers.hpp b/cpp/src/jit/common_headers.hpp index 430f44e2fb8..876ecce5194 100644 --- a/cpp/src/jit/common_headers.hpp +++ b/cpp/src/jit/common_headers.hpp @@ -21,14 +21,18 @@ #include #include #include +#include #include #include +#include #include #include #include #include +#include #include #include +#include #include #include #include @@ -69,6 +73,8 @@ const std::unordered_map stringified_headers{ {"simt/../../libcxx/include/ctime", libcxx_ctime}, {"simt/../../libcxx/include/limits", libcxx_limits}, {"simt/../../libcxx/include/ratio", libcxx_ratio}, + {"simt/../../libcxx/include/cmath", libcxx_cmath}, + {"simt/../../libcxx/include/cassert", libcxx_cassert}, {"simt/../../libcxx/include/type_traits", libcxx_type_traits}, {"simt/../details/__config", libcudacxx_details_config}, {"simt/cfloat", libcudacxx_simt_cfloat}, diff --git a/cpp/src/jit/type.cpp b/cpp/src/jit/type.cpp index 8f36f9a7880..60ebc97e20d 100644 --- a/cpp/src/jit/type.cpp +++ b/cpp/src/jit/type.cpp @@ -63,8 +63,7 @@ const void* get_data_ptr(column_view const& view) const void* get_data_ptr(scalar const& s) { - auto val = type_dispatcher(s.type(), get_data_ptr_functor{}, s); - return val; + return type_dispatcher(s.type(), get_data_ptr_functor{}, s); } std::string get_type_name(data_type type) diff --git a/cpp/src/reductions/simple.cuh b/cpp/src/reductions/simple.cuh index cf5a30b1c46..f39ad0b27b8 100644 --- a/cpp/src/reductions/simple.cuh +++ b/cpp/src/reductions/simple.cuh @@ -76,10 +76,12 @@ struct result_type_dispatcher { // - same dtypes (including cudf::wrappers) // - any arithmetic dtype to any arithmetic dtype // - bool to/from any arithmetic dtype + // - fixed_point to fixed_point return cudf::is_convertible::value && (std::is_arithmetic::value || std::is_same::value || - std::is_same::value) && + std::is_same::value || + cudf::is_fixed_point()) && !std::is_same::value; } diff --git a/cpp/src/replace/clamp.cu b/cpp/src/replace/clamp.cu index d9b2c50c081..a47892baf39 100644 --- a/cpp/src/replace/clamp.cu +++ b/cpp/src/replace/clamp.cu @@ -286,6 +286,32 @@ std::unique_ptr dispatch_clamp::operator()( CUDF_FAIL("clamp for list_view not supported"); } +template <> +std::unique_ptr dispatch_clamp::operator()( + column_view const& input, + scalar const& lo, + scalar const& lo_replace, + scalar const& hi, + scalar const& hi_replace, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) +{ + CUDF_FAIL("clamp for decimal32 not supported"); +} + +template <> +std::unique_ptr dispatch_clamp::operator()( + column_view const& input, + scalar const& lo, + scalar const& lo_replace, + scalar const& hi, + scalar const& hi_replace, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) +{ + CUDF_FAIL("clamp for decimal32 not supported"); +} + /** * @copydoc cudf::clamp(column_view const& input, scalar const& lo, diff --git a/cpp/src/replace/replace.cu b/cpp/src/replace/replace.cu index 142ffb5a778..f18ce11b937 100644 --- a/cpp/src/replace/replace.cu +++ b/cpp/src/replace/replace.cu @@ -75,7 +75,7 @@ __device__ auto get_new_value(cudf::size_type idx, { auto found_ptr = thrust::find(thrust::seq, values_to_replace_begin, values_to_replace_end, input_data[idx]); - T new_value{0}; + T new_value{}; bool output_is_valid{true}; if (found_ptr != values_to_replace_end) { diff --git a/cpp/src/scalar/scalar_factories.cpp b/cpp/src/scalar/scalar_factories.cpp index aa0088c4dc0..6bff618c815 100644 --- a/cpp/src/scalar/scalar_factories.cpp +++ b/cpp/src/scalar/scalar_factories.cpp @@ -27,7 +27,7 @@ struct scalar_construction_helper { std::enable_if_t(), std::unique_ptr> operator()( cudaStream_t stream, rmm::mr::device_memory_resource* mr) const { - auto s = new ScalarType(0, false, stream, mr); + auto s = new ScalarType(T{}, false, stream, mr); return std::unique_ptr(s); } diff --git a/cpp/src/unary/cast_ops.cu b/cpp/src/unary/cast_ops.cu index 3051e1f8980..338bb481606 100644 --- a/cpp/src/unary/cast_ops.cu +++ b/cpp/src/unary/cast_ops.cu @@ -103,7 +103,7 @@ struct dispatch_unary_cast_to { // Disallow conversions between timestamps and numeric template < typename TargetT, - typename std::enable_if_t() && + typename std::enable_if_t() && !cudf::is_fixed_point() && !(cudf::is_timestamp() && is_numeric()) && !(cudf::is_timestamp() && is_numeric())>* = nullptr> std::unique_ptr operator()(data_type type, @@ -130,7 +130,7 @@ struct dispatch_unary_cast_to { template < typename TargetT, - typename std::enable_if_t() || + typename std::enable_if_t() || cudf::is_fixed_point() || (cudf::is_timestamp() && is_numeric()) || (cudf::is_timestamp() && is_numeric())>* = nullptr> std::unique_ptr operator()(data_type type, @@ -139,6 +139,8 @@ struct dispatch_unary_cast_to { { if (!cudf::is_fixed_width()) CUDF_FAIL("Column type must be numeric or chrono"); + else if (cudf::is_fixed_point()) + CUDF_FAIL("Fixed point unary ops currently not supported"); else if (cudf::is_timestamp() && is_numeric()) CUDF_FAIL("Timestamps can be created only from duration"); else @@ -151,7 +153,9 @@ struct dispatch_unary_cast_from { dispatch_unary_cast_from(column_view inp) : input(inp) {} - template ()>* = nullptr> + template < + typename T, + typename std::enable_if_t() && !cudf::is_fixed_point()>* = nullptr> std::unique_ptr operator()(data_type type, rmm::mr::device_memory_resource* mr, cudaStream_t stream) @@ -159,6 +163,14 @@ struct dispatch_unary_cast_from { return type_dispatcher(type, dispatch_unary_cast_to{input}, type, mr, stream); } + template ()>* = nullptr> + std::unique_ptr operator()(data_type type, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) + { + CUDF_FAIL("Fixed point unary ops not supported yet"); + } + template ()>* = nullptr> std::unique_ptr operator()(data_type type, rmm::mr::device_memory_resource* mr, diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 75d5f00d38d..e40c6b02504 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -18,8 +18,14 @@ */ #include -#include #include +#include + +#include +#include +#include "cudf/types.hpp" +#include "cudf/utilities/type_dispatcher.hpp" +#include "tests/utilities/column_utilities.hpp" namespace cudf { namespace test { @@ -2036,6 +2042,141 @@ TEST_F(BinaryOperationIntegrationTest, ATan2_Vector_Vector_FP64_SI32_SI64) ASSERT_BINOP(*out, lhs, rhs, ATAN2(), NearEqualComparator{2}); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +template +using wrapper = cudf::test::fixed_width_column_wrapper; +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto const sz = std::size_t{1000}; + + auto vec1 = std::vector(sz); + auto const vec2 = std::vector(sz, decimalXX{1, scale_type{-1}}); + auto expected = std::vector(sz); + + std::iota(std::begin(vec1), std::end(vec1), decimalXX{}); + + std::transform(std::cbegin(vec1), + std::cend(vec1), + std::cbegin(vec2), + std::begin(expected), + std::plus()); + + auto const lhs = wrapper(vec1.begin(), vec1.end()); + auto const rhs = wrapper(vec2.begin(), vec2.end()); + auto const expected_col = wrapper(expected.begin(), expected.end()); + + auto const result = cudf::binary_operation( + lhs, rhs, cudf::binary_operator::ADD, static_cast(lhs).type()); + + cudf::test::expect_columns_equal(expected_col, result->view()); +} + +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiply) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto const sz = std::size_t{1000}; + + auto vec1 = std::vector(sz); + auto const vec2 = std::vector(sz, decimalXX{1, scale_type{-1}}); + auto expected = std::vector(sz); + + std::iota(std::begin(vec1), std::end(vec1), decimalXX{}); + + std::transform(std::cbegin(vec1), + std::cend(vec1), + std::cbegin(vec2), + std::begin(expected), + std::multiplies()); + + auto const lhs = wrapper(vec1.begin(), vec1.end()); + auto const rhs = wrapper(vec2.begin(), vec2.end()); + auto const expected_col = wrapper(expected.begin(), expected.end()); + + auto const result = cudf::binary_operation( + lhs, rhs, cudf::binary_operator::MUL, static_cast(lhs).type()); + + cudf::test::expect_columns_equal(expected_col, result->view()); +} + +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimple) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto const ONE = decimalXX{1, scale_type{0}}; + auto const TWO = decimalXX{2, scale_type{0}}; + auto const THREE = decimalXX{3, scale_type{0}}; + auto const FOUR = decimalXX{4, scale_type{0}}; + + auto const ONE_2 = decimalXX{1, scale_type{-2}}; + auto const TWO_2 = decimalXX{2, scale_type{-2}}; + auto const THREE_2 = decimalXX{3, scale_type{-2}}; + auto const FOUR_2 = decimalXX{4, scale_type{-2}}; + + auto const vec1 = std::vector{ONE, TWO, THREE, FOUR}; + auto const vec2 = std::vector{ONE_2, TWO_2, THREE_2, FOUR_2}; + auto const trues = std::vector(4, true); + + auto const col1 = wrapper(vec1.begin(), vec1.end()); + auto const col2 = wrapper(vec2.begin(), vec2.end()); + auto const expected = wrapper(trues.begin(), trues.end()); + + auto const result = cudf::binary_operation( + col1, col2, cudf::binary_operator::EQUAL, cudf::data_type{type_id::BOOL8}); + + cudf::test::expect_columns_equal(expected, result->view()); +} + +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualLessGreater) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto const sz = std::size_t{1000}; + + // TESTING binary op ADD + + auto vec1 = std::vector(sz, decimalXX{0, scale_type{-3}}); + auto vec2 = std::vector(sz, decimalXX{0, scale_type{-1}}); + + std::iota(std::begin(vec1), std::end(vec1), decimalXX{1, scale_type{-3}}); + + auto const iota_1 = wrapper(vec1.begin(), vec1.end()); + auto const zeros_3 = wrapper(vec2.begin(), vec2.end()); + + auto const iota_3 = cudf::binary_operation( + zeros_3, iota_1, cudf::binary_operator::ADD, static_cast(zeros_3).type()); + + cudf::test::expect_columns_equal(iota_1, iota_3->view()); + + // TESTING binary op EQUAL, LESS, GREATER + + auto const trues = std::vector(sz, true); + auto const true_col = wrapper(trues.begin(), trues.end()); + + auto const equal_result = cudf::binary_operation( + iota_1, iota_3->view(), cudf::binary_operator::EQUAL, data_type{type_id::BOOL8}); + cudf::test::expect_columns_equal(true_col, equal_result->view()); + + auto const less_result = cudf::binary_operation( + zeros_3, iota_3->view(), cudf::binary_operator::LESS, data_type{type_id::BOOL8}); + cudf::test::expect_columns_equal(true_col, less_result->view()); + + auto const greater_result = cudf::binary_operation( + iota_3->view(), zeros_3, cudf::binary_operator::GREATER, data_type{type_id::BOOL8}); + cudf::test::expect_columns_equal(true_col, greater_result->view()); +} + } // namespace binop } // namespace test } // namespace cudf diff --git a/cpp/tests/column/factories_test.cpp b/cpp/tests/column/factories_test.cpp index bfb16da3871..7ae845c13db 100644 --- a/cpp/tests/column/factories_test.cpp +++ b/cpp/tests/column/factories_test.cpp @@ -199,7 +199,7 @@ template class FixedWidthFactoryTest : public ColumnFactoryTest { }; -TYPED_TEST_CASE(FixedWidthFactoryTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(FixedWidthFactoryTest, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(FixedWidthFactoryTest, EmptyNoMask) { diff --git a/cpp/tests/copying/concatenate_tests.cu b/cpp/tests/copying/concatenate_tests.cu index e577857723c..449c6c4558a 100644 --- a/cpp/tests/copying/concatenate_tests.cu +++ b/cpp/tests/copying/concatenate_tests.cu @@ -13,10 +13,6 @@ * See the License for the specific language governing permissions and * limitations under the License. */ -#include -#include -#include -#include #include #include @@ -24,6 +20,12 @@ #include #include +#include +#include +#include +#include +#include + #include template @@ -546,3 +548,30 @@ TEST_F(ListsColumnTest, ConcatenateIncompleteHierarchies) cudf::test::expect_columns_equal(*result, expected); } } + +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +template +using wrapper = cudf::test::fixed_width_column_wrapper; +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointConcatentate) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto vec = std::vector(1000); + std::iota(std::begin(vec), std::end(vec), decimalXX{}); + + auto const a = wrapper(vec.begin(), /***/ vec.begin() + 300); + auto const b = wrapper(vec.begin() + 300, vec.begin() + 700); + auto const c = wrapper(vec.begin() + 700, vec.end()); + + auto const fixed_point_columns = std::vector{a, b, c}; + auto const results = cudf::concatenate(fixed_point_columns); + auto const expected = wrapper(vec.begin(), vec.end()); + + cudf::test::expect_columns_equal(*results, expected); +} \ No newline at end of file diff --git a/cpp/tests/copying/copy_tests.cu b/cpp/tests/copying/copy_tests.cu index bc2d8c76f34..3e3685fe1d0 100644 --- a/cpp/tests/copying/copy_tests.cu +++ b/cpp/tests/copying/copy_tests.cu @@ -32,7 +32,7 @@ template struct CopyTest : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(CopyTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(CopyTest, cudf::test::FixedWidthTypesWithoutFixedPoint); #define wrapper cudf::test::fixed_width_column_wrapper diff --git a/cpp/tests/copying/get_value_tests.cpp b/cpp/tests/copying/get_value_tests.cpp index e6c4ae6fe7c..742296cdaf3 100644 --- a/cpp/tests/copying/get_value_tests.cpp +++ b/cpp/tests/copying/get_value_tests.cpp @@ -33,7 +33,7 @@ template struct FixedWidthGetValueTest : public BaseFixture { }; -TYPED_TEST_CASE(FixedWidthGetValueTest, FixedWidthTypes); +TYPED_TEST_CASE(FixedWidthGetValueTest, FixedWidthTypesWithoutFixedPoint); TYPED_TEST(FixedWidthGetValueTest, BasicGet) { @@ -123,7 +123,7 @@ template struct DictionaryGetValueTest : public BaseFixture { }; -TYPED_TEST_CASE(DictionaryGetValueTest, FixedWidthTypes); +TYPED_TEST_CASE(DictionaryGetValueTest, FixedWidthTypesWithoutFixedPoint); TYPED_TEST(DictionaryGetValueTest, BasicGet) { diff --git a/cpp/tests/copying/scatter_tests.cpp b/cpp/tests/copying/scatter_tests.cpp index 3990577e264..f6ab4817618 100644 --- a/cpp/tests/copying/scatter_tests.cpp +++ b/cpp/tests/copying/scatter_tests.cpp @@ -255,8 +255,9 @@ class ScatterInvalidIndexTypeTests : public cudf::test::BaseFixture { }; // NOTE string types hit static assert in fixed_width_column_wrapper -using InvalidIndexTypes = - cudf::test::Concat, cudf::test::ChronoTypes>; +using InvalidIndexTypes = cudf::test::Concat, + cudf::test::ChronoTypes, + cudf::test::FixedPointTypes>; TYPED_TEST_CASE(ScatterInvalidIndexTypeTests, InvalidIndexTypes); // Throw logic error if scatter map column has invalid data type @@ -301,7 +302,7 @@ template class ScatterDataTypeTests : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(ScatterDataTypeTests, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(ScatterDataTypeTests, cudf::test::FixedWidthTypesWithoutFixedPoint); // Empty scatter map returns copy of input TYPED_TEST(ScatterDataTypeTests, EmptyScatterMap) @@ -795,7 +796,7 @@ struct BooleanMaskScalarScatter : public cudf::test::BaseFixture { } }; -TYPED_TEST_CASE(BooleanMaskScalarScatter, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(BooleanMaskScalarScatter, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(BooleanMaskScalarScatter, WithNoNullElementsInTarget) { @@ -948,3 +949,36 @@ TEST_F(BooleanMaskScatterScalarFails, NumberOfColumnAndScalarMismatch) EXPECT_THROW(cudf::boolean_mask_scatter(scalar_vect, target_table, mask), cudf::logic_error); } + +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +template +using wrapper = cudf::test::fixed_width_column_wrapper; +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointScatter) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto const ONE = decimalXX{1, scale_type{0}}; + auto const TWO = decimalXX{2, scale_type{0}}; + auto const THREE = decimalXX{3, scale_type{0}}; + auto const FOUR = decimalXX{4, scale_type{0}}; + auto const FIVE = decimalXX{5, scale_type{0}}; + + auto const source = wrapper({ONE, TWO, THREE, FOUR, FIVE}); + auto const target = wrapper({ONE, TWO, THREE, FOUR, FIVE, FOUR, THREE, TWO, ONE}); + auto const scatter_map = wrapper({1, 2, -1, -3, -4}); + auto const expected = wrapper({ONE, ONE, TWO, FOUR, FIVE, FIVE, FOUR, TWO, THREE}); + + auto const source_table = cudf::table_view({source, source}); + auto const target_table = cudf::table_view({target, target}); + auto const expected_table = cudf::table_view({expected, expected}); + + auto const result = cudf::scatter(source_table, scatter_map, target_table, true); + + cudf::test::expect_tables_equal(expected_table, result->view()); +} diff --git a/cpp/tests/copying/shift_tests.cpp b/cpp/tests/copying/shift_tests.cpp index 2d864438d6b..3cc634558df 100644 --- a/cpp/tests/copying/shift_tests.cpp +++ b/cpp/tests/copying/shift_tests.cpp @@ -34,7 +34,7 @@ template > std::unique_ptr make_scalar( cudaStream_t stream = 0, rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource()) { - auto s = new ScalarType(T{0}, false, stream, mr); + auto s = new ScalarType(T{}, false, stream, mr); return std::unique_ptr(s); } @@ -58,7 +58,7 @@ template struct ShiftTest : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(ShiftTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(ShiftTest, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(ShiftTest, OneColumnEmpty) { diff --git a/cpp/tests/device_atomics/device_atomics_test.cu b/cpp/tests/device_atomics/device_atomics_test.cu index 6922ea75690..b94e4498d50 100644 --- a/cpp/tests/device_atomics/device_atomics_test.cu +++ b/cpp/tests/device_atomics/device_atomics_test.cu @@ -162,7 +162,7 @@ struct AtomicsTest : public cudf::test::BaseFixture { } }; -TYPED_TEST_CASE(AtomicsTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(AtomicsTest, cudf::test::FixedWidthTypesWithoutFixedPoint); // tests for atomicAdd/Min/Max TYPED_TEST(AtomicsTest, atomicOps) diff --git a/cpp/tests/filling/fill_tests.cpp b/cpp/tests/filling/fill_tests.cpp index 4802fb1416d..8b7b541132b 100644 --- a/cpp/tests/filling/fill_tests.cpp +++ b/cpp/tests/filling/fill_tests.cpp @@ -93,7 +93,7 @@ class FillTypedTestFixture : public cudf::test::BaseFixture { } }; -TYPED_TEST_CASE(FillTypedTestFixture, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(FillTypedTestFixture, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(FillTypedTestFixture, SetSingle) { diff --git a/cpp/tests/filling/repeat_tests.cpp b/cpp/tests/filling/repeat_tests.cpp index f585df527e6..529d6b419a2 100644 --- a/cpp/tests/filling/repeat_tests.cpp +++ b/cpp/tests/filling/repeat_tests.cpp @@ -42,7 +42,7 @@ class RepeatTypedTestFixture : public cudf::test::BaseFixture, cudf::size_type repeat_count() { return this->generate(); } }; -TYPED_TEST_CASE(RepeatTypedTestFixture, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(RepeatTypedTestFixture, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(RepeatTypedTestFixture, RepeatScalarCount) { diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 29cf87a8e13..cfd669f8246 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -14,14 +14,16 @@ * limitations under the License. */ -#include +#include +#include + #include #include #include + +#include #include #include -#include -#include #include #include diff --git a/cpp/tests/groupby/group_max_test.cpp b/cpp/tests/groupby/group_max_test.cpp index c02057478a2..5acdd5546e6 100644 --- a/cpp/tests/groupby/group_max_test.cpp +++ b/cpp/tests/groupby/group_max_test.cpp @@ -28,7 +28,7 @@ template struct groupby_max_test : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(groupby_max_test, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(groupby_max_test, cudf::test::FixedWidthTypesWithoutFixedPoint); // clang-format off TYPED_TEST(groupby_max_test, basic) diff --git a/cpp/tests/groupby/group_min_test.cpp b/cpp/tests/groupby/group_min_test.cpp index 3240d228079..7c7ae948f3c 100644 --- a/cpp/tests/groupby/group_min_test.cpp +++ b/cpp/tests/groupby/group_min_test.cpp @@ -28,7 +28,7 @@ template struct groupby_min_test : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(groupby_min_test, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(groupby_min_test, cudf::test::FixedWidthTypesWithoutFixedPoint); // clang-format off TYPED_TEST(groupby_min_test, basic) diff --git a/cpp/tests/grouped_rolling/grouped_rolling_test.cpp b/cpp/tests/grouped_rolling/grouped_rolling_test.cpp index 460a4c3ad15..6dfa74b644f 100644 --- a/cpp/tests/grouped_rolling/grouped_rolling_test.cpp +++ b/cpp/tests/grouped_rolling/grouped_rolling_test.cpp @@ -536,7 +536,7 @@ TEST_F(GroupedRollingErrorTest, SumTimestampNotSupported) cudf::logic_error); } -TYPED_TEST_CASE(GroupedRollingTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(GroupedRollingTest, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(GroupedRollingTest, SimplePartitionedStaticWindowsWithGroupKeys) { diff --git a/cpp/tests/merge/merge_string_test.cpp b/cpp/tests/merge/merge_string_test.cpp index 9cafcf02ab0..e077a1cfc22 100644 --- a/cpp/tests/merge/merge_string_test.cpp +++ b/cpp/tests/merge/merge_string_test.cpp @@ -44,7 +44,7 @@ template class MergeStringTest : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(MergeStringTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(MergeStringTest, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(MergeStringTest, Merge1StringKeyColumns) { diff --git a/cpp/tests/merge/merge_test.cpp b/cpp/tests/merge/merge_test.cpp index 4926ebbff80..f848fc58ae2 100644 --- a/cpp/tests/merge/merge_test.cpp +++ b/cpp/tests/merge/merge_test.cpp @@ -42,7 +42,7 @@ template class MergeTest_ : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(MergeTest_, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(MergeTest_, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(MergeTest_, MergeIsZeroWhenShouldNotBeZero) { diff --git a/cpp/tests/partitioning/hash_partition_test.cpp b/cpp/tests/partitioning/hash_partition_test.cpp index ffd9d7fe7ed..e8ce9a040da 100644 --- a/cpp/tests/partitioning/hash_partition_test.cpp +++ b/cpp/tests/partitioning/hash_partition_test.cpp @@ -186,7 +186,7 @@ template class HashPartitionFixedWidth : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(HashPartitionFixedWidth, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(HashPartitionFixedWidth, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(HashPartitionFixedWidth, NullableFixedWidth) { diff --git a/cpp/tests/partitioning/round_robin_test.cpp b/cpp/tests/partitioning/round_robin_test.cpp index 3ee748ef111..e0c42b64fc5 100644 --- a/cpp/tests/partitioning/round_robin_test.cpp +++ b/cpp/tests/partitioning/round_robin_test.cpp @@ -45,7 +45,7 @@ template class RoundRobinTest : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(RoundRobinTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(RoundRobinTest, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(RoundRobinTest, RoundRobinPartitions13_3) { diff --git a/cpp/tests/reductions/reduction_tests.cpp b/cpp/tests/reductions/reduction_tests.cpp index 999f4f8f6ca..520b1890632 100644 --- a/cpp/tests/reductions/reduction_tests.cpp +++ b/cpp/tests/reductions/reduction_tests.cpp @@ -22,6 +22,7 @@ #include #include +#include #include #include @@ -920,4 +921,62 @@ TYPED_TEST(ReductionTest, UniqueCount) cudf::make_nunique_aggregation(cudf::null_policy::EXCLUDE)); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointReductionProduct) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto const ONE = decimalXX{1, scale_type{0}}; + auto const TWO = decimalXX{2, scale_type{0}}; + auto const THREE = decimalXX{3, scale_type{0}}; + auto const FOUR = decimalXX{4, scale_type{0}}; + // auto const _24 = decimalXX{24, scale_type{0}}; + + auto const in = std::vector{ONE, TWO, THREE, FOUR}; + auto const column = cudf::test::fixed_width_column_wrapper(in.cbegin(), in.cend()); + // auto const expected = std::accumulate(in.cbegin(), in.cend(), ONE, + // std::multiplies()); + auto const out_type = static_cast(column).type(); + + EXPECT_THROW(cudf::reduce(column, cudf::make_product_aggregation(), out_type), cudf::logic_error); + + // auto const result = cudf::reduce(column, cudf::make_product_aggregation(), out_type); + // auto const result_scalar = static_cast*>(result.get()); + + // EXPECT_EQ(result_scalar->value(), expected); + // EXPECT_EQ(result_scalar->value(), _24); +} + +TYPED_TEST(FixedPointTestBothReps, FixedPointReductionSum) +{ + using namespace numeric; + using decimalXX = TypeParam; + + // auto const ZERO = decimalXX{0, scale_type{0}}; + auto const ONE = decimalXX{1, scale_type{0}}; + auto const TWO = decimalXX{2, scale_type{0}}; + auto const THREE = decimalXX{3, scale_type{0}}; + auto const FOUR = decimalXX{4, scale_type{0}}; + // auto const TEN = decimalXX{10, scale_type{0}}; + + auto const in = std::vector{ONE, TWO, THREE, FOUR}; + auto const column = cudf::test::fixed_width_column_wrapper(in.cbegin(), in.cend()); + // auto const expected = std::accumulate(in.cbegin(), in.cend(), ZERO, std::plus()); + auto const out_type = static_cast(column).type(); + + EXPECT_THROW(cudf::reduce(column, cudf::make_sum_aggregation(), out_type), cudf::logic_error); + + // auto const result = cudf::reduce(column, cudf::make_sum_aggregation(), out_type); + // auto const result_scalar = static_cast*>(result.get()); + + // EXPECT_EQ(result_scalar->value(), expected); + // EXPECT_EQ(result_scalar->value(), TEN); +} + CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/replace/clamp_test.cpp b/cpp/tests/replace/clamp_test.cpp index 210df46cf45..735a66011ca 100644 --- a/cpp/tests/replace/clamp_test.cpp +++ b/cpp/tests/replace/clamp_test.cpp @@ -200,7 +200,7 @@ struct ClampTestNumeric : public cudf::test::BaseFixture { } } }; -using Types = cudf::test::FixedWidthTypes; +using Types = cudf::test::FixedWidthTypesWithoutFixedPoint; TYPED_TEST_CASE(ClampTestNumeric, Types); diff --git a/cpp/tests/replace/replace_tests.cpp b/cpp/tests/replace/replace_tests.cpp index ef0ac5dcebf..fd3352d6c37 100644 --- a/cpp/tests/replace/replace_tests.cpp +++ b/cpp/tests/replace/replace_tests.cpp @@ -16,14 +16,16 @@ * See the License for the specific language governing permissions and * limitations under the License. */ -#include -#include #include #include #include #include #include +#include +#include +#include "cudf/fixed_point/fixed_point.hpp" + #include #include @@ -528,4 +530,39 @@ TYPED_TEST(ReplaceTest, LargeScaleReplaceTest) expect_columns_equal(expected, *actual_result); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +template +using wrapper = cudf::test::fixed_width_column_wrapper; +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointReplace) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto const ONE = decimalXX{1, scale_type{0}}; + auto const TWO = decimalXX{2, scale_type{0}}; + auto const sz = std::size_t{1000}; + + auto vec1 = std::vector(sz); + auto const vec2 = std::vector(sz, TWO); + + std::generate(vec1.begin(), vec1.end(), [&, i = 0]() mutable { return ++i % 2 ? ONE : TWO; }); + + auto const to_replace = std::vector{ONE}; + auto const replacement = std::vector{TWO}; + + auto const input_w = wrapper(vec1.begin(), vec1.end()); + auto const to_replace_w = wrapper(to_replace.begin(), to_replace.end()); + auto const replacement_w = wrapper(replacement.begin(), replacement.end()); + auto const expected_w = wrapper(vec2.begin(), vec2.end()); + + auto const result = cudf::find_and_replace_all(input_w, to_replace_w, replacement_w); + + cudf::test::expect_columns_equal(*result, expected_w); +} + CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index bdf377a1dd1..4497b0d3e52 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -21,7 +21,9 @@ #include #include +#include #include + #include using namespace cudf::test; @@ -60,10 +62,7 @@ TYPED_TEST(InterleaveColumnsTest, TwoColumns) auto a = fixed_width_column_wrapper({0, 2}); auto b = fixed_width_column_wrapper({1, 3}); - cudf::table_view in(std::vector{ - a, - b, - }); + cudf::table_view in(std::vector{a, b}); auto expected = fixed_width_column_wrapper({0, 1, 2, 3}); auto actual = cudf::interleave_columns(in); @@ -346,4 +345,32 @@ TEST_F(InterleaveStringsColumnsTest, MultiColumnStringMixNullableMix) cudf::test::expect_columns_equal(*results, exp_results, true); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointInterleave) +{ + using namespace numeric; + using decimalXX = TypeParam; + + for (int i = 0; i > -4; --i) { + auto const ONE = decimalXX{1, scale_type{i}}; + auto const TWO = decimalXX{2, scale_type{i}}; + auto const FOUR = decimalXX{4, scale_type{i}}; + auto const FIVE = decimalXX{5, scale_type{i}}; + + auto const a = cudf::test::fixed_width_column_wrapper({ONE, FOUR}); + auto const b = cudf::test::fixed_width_column_wrapper({TWO, FIVE}); + + auto const input = cudf::table_view{std::vector{a, b}}; + auto const expected = cudf::test::fixed_width_column_wrapper({ONE, TWO, FOUR, FIVE}); + auto const actual = cudf::interleave_columns(input); + + expect_columns_equal(expected, actual->view()); + } +} + CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/rolling/rolling_test.cpp b/cpp/tests/rolling/rolling_test.cpp index 1ee7c663f4e..103b859195c 100644 --- a/cpp/tests/rolling/rolling_test.cpp +++ b/cpp/tests/rolling/rolling_test.cpp @@ -523,7 +523,7 @@ TEST_F(RollingErrorTest, MeanTimestampNotSupported) cudf::logic_error); } -TYPED_TEST_CASE(RollingTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(RollingTest, cudf::test::FixedWidthTypesWithoutFixedPoint); // simple example from Pandas docs TYPED_TEST(RollingTest, SimpleStatic) diff --git a/cpp/tests/scalar/factories_test.cpp b/cpp/tests/scalar/factories_test.cpp index 7307aaabeed..24a88268ab7 100644 --- a/cpp/tests/scalar/factories_test.cpp +++ b/cpp/tests/scalar/factories_test.cpp @@ -118,7 +118,7 @@ template struct FixedWidthScalarFactory : public ScalarFactoryTest { }; -TYPED_TEST_CASE(FixedWidthScalarFactory, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(FixedWidthScalarFactory, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(FixedWidthScalarFactory, ValueProvided) { diff --git a/cpp/tests/scalar/scalar_device_view_test.cu b/cpp/tests/scalar/scalar_device_view_test.cu index 2e842ffc5dc..e86f6a0bc2a 100644 --- a/cpp/tests/scalar/scalar_device_view_test.cu +++ b/cpp/tests/scalar/scalar_device_view_test.cu @@ -30,7 +30,7 @@ template struct TypedScalarDeviceViewTest : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(TypedScalarDeviceViewTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(TypedScalarDeviceViewTest, cudf::test::FixedWidthTypesWithoutFixedPoint); template __global__ void test_set_value(ScalarDeviceViewType s, ScalarDeviceViewType s1) diff --git a/cpp/tests/scalar/scalar_test.cpp b/cpp/tests/scalar/scalar_test.cpp index 4464dd0ea00..bac15542e75 100644 --- a/cpp/tests/scalar/scalar_test.cpp +++ b/cpp/tests/scalar/scalar_test.cpp @@ -29,7 +29,7 @@ template struct TypedScalarTest : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(TypedScalarTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(TypedScalarTest, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(TypedScalarTest, DefaultValidity) { diff --git a/cpp/tests/search/search_test.cpp b/cpp/tests/search/search_test.cpp index 6e2ace07b99..4ca8de65e81 100644 --- a/cpp/tests/search/search_test.cpp +++ b/cpp/tests/search/search_test.cpp @@ -17,7 +17,9 @@ #include #include #include +#include +#include #include "cudf/search.hpp" struct SearchTest : public cudf::test::BaseFixture { @@ -1815,4 +1817,58 @@ TEST_F(SearchTest, multi_contains_empty_input_set_string) expect_columns_equal(*result, expect); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointLowerBound) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto vec = std::vector(1000); + std::iota(std::begin(vec), std::end(vec), decimalXX{}); + + auto const values = + cudf::test::fixed_width_column_wrapper{decimalXX{200, scale_type{0}}, + decimalXX{400, scale_type{0}}, + decimalXX{600, scale_type{0}}, + decimalXX{800, scale_type{0}}}; + auto const expect = cudf::test::fixed_width_column_wrapper{200, 400, 600, 800}; + auto const column = cudf::test::fixed_width_column_wrapper(vec.begin(), vec.end()); + + auto result = cudf::lower_bound({cudf::table_view{{column}}}, + {cudf::table_view{{values}}}, + {cudf::order::ASCENDING}, + {cudf::null_order::BEFORE}); + + expect_columns_equal(*result, expect); +} + +TYPED_TEST(FixedPointTestBothReps, FixedPointUpperBound) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto vec = std::vector(1000); + std::iota(std::begin(vec), std::end(vec), decimalXX{}); + + auto const values = + cudf::test::fixed_width_column_wrapper{decimalXX{200, scale_type{0}}, + decimalXX{400, scale_type{0}}, + decimalXX{600, scale_type{0}}, + decimalXX{800, scale_type{0}}}; + auto const expect = cudf::test::fixed_width_column_wrapper{201, 401, 601, 801}; + auto const column = cudf::test::fixed_width_column_wrapper(vec.begin(), vec.end()); + + auto result = cudf::upper_bound({cudf::table_view{{column}}}, + {cudf::table_view{{values}}}, + {cudf::order::ASCENDING}, + {cudf::null_order::BEFORE}); + + expect_columns_equal(*result, expect); +} + CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/sort/sort_test.cpp b/cpp/tests/sort/sort_test.cpp index bc4737c4cb1..0d68047d0d4 100644 --- a/cpp/tests/sort/sort_test.cpp +++ b/cpp/tests/sort/sort_test.cpp @@ -14,17 +14,20 @@ * limitations under the License. */ +#include +#include +#include +#include +#include + #include #include +#include #include #include #include #include -#include -#include -#include -#include -#include + #include namespace cudf { @@ -282,6 +285,43 @@ TEST_F(SortByKey, ValueKeysSizeMismatch) EXPECT_THROW(sort_by_key(values, keys), logic_error); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +template +using wrapper = cudf::test::fixed_width_column_wrapper; +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointSortedOrderGather) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto const ZERO = decimalXX{0, scale_type{0}}; + auto const ONE = decimalXX{1, scale_type{0}}; + auto const TWO = decimalXX{2, scale_type{0}}; + auto const THREE = decimalXX{3, scale_type{0}}; + auto const FOUR = decimalXX{4, scale_type{0}}; + + auto const input_vec = std::vector{TWO, ONE, ZERO, FOUR, THREE}; + auto const index_vec = std::vector{2, 1, 0, 4, 3}; + auto const sorted_vec = std::vector{ZERO, ONE, TWO, THREE, FOUR}; + + auto const input_col = wrapper(input_vec.begin(), input_vec.end()); + auto const index_col = wrapper(index_vec.begin(), index_vec.end()); + auto const sorted_col = wrapper(sorted_vec.begin(), sorted_vec.end()); + + auto const sorted_table = cudf::table_view{{sorted_col}}; + auto const input_table = cudf::table_view{{input_col}}; + + auto const indices = cudf::sorted_order(input_table); + auto const sorted = cudf::gather(input_table, indices->view()); + + cudf::test::expect_columns_equal(index_col, indices->view()); + cudf::test::expect_tables_equal(sorted_table, sorted->view()); +} + } // namespace test } // namespace cudf diff --git a/cpp/tests/transpose/transpose_test.cpp b/cpp/tests/transpose/transpose_test.cpp index 8730b985e32..1d86ed9ff20 100644 --- a/cpp/tests/transpose/transpose_test.cpp +++ b/cpp/tests/transpose/transpose_test.cpp @@ -153,7 +153,7 @@ template class TransposeTest : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(TransposeTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(TransposeTest, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(TransposeTest, SingleValue) { run_test(1, 1, false); } diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index 636daf07ab7..1d985c7b6ed 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -440,6 +440,23 @@ struct column_view_printer { this->template operator()(*col_as_strings, out, indent); } + template ()>* = nullptr> + void operator()(cudf::column_view const& col, + std::vector& out, + std::string const& indent) + { + auto const h_data = cudf::test::to_host(col); + + out.resize(col.size()); + std::transform(thrust::make_counting_iterator(size_type{0}), + thrust::make_counting_iterator(col.size()), + out.begin(), + [&](auto idx) { + auto const d = static_cast(h_data.first[idx]); + return std::to_string(d); + }); + } + template ::value>* = nullptr> void operator()(cudf::column_view const& col, diff --git a/cpp/tests/utilities/column_wrapper.hpp b/cpp/tests/utilities/column_wrapper.hpp index bfd8a387d67..6238c833e7b 100644 --- a/cpp/tests/utilities/column_wrapper.hpp +++ b/cpp/tests/utilities/column_wrapper.hpp @@ -35,6 +35,7 @@ #include #include +#include "cudf/fixed_point/fixed_point.hpp" namespace cudf { namespace test { @@ -132,6 +133,18 @@ struct fixed_width_type_converter { std::transform(begin, end, out, [](auto const& e) { return static_cast(e); }); } + // Is the target type a fixed_point type + template ::value && + cudf::is_fixed_point()>::type* = nullptr> + void operator()(InputIterator begin, InputIterator end, OutputIterator out) const + { + std::transform(begin, end, out, [](auto const& e) { return ToT{e, numeric::scale_type{0}}; }); + } + #if 0 // This is to be used when timestamp disallows construction from tick counts; presently, // this conflicts with the convertible/constructible overload diff --git a/cpp/tests/utilities/type_lists.hpp b/cpp/tests/utilities/type_lists.hpp index 60dfbe5a3c9..00d41e835fb 100644 --- a/cpp/tests/utilities/type_lists.hpp +++ b/cpp/tests/utilities/type_lists.hpp @@ -24,6 +24,7 @@ #include #include +#include "cudf/fixed_point/fixed_point.hpp" /** * @filename type_lists.hpp @@ -182,6 +183,18 @@ using StringTypes = cudf::test::Types; */ using ListTypes = cudf::test::Types; +/** + * @brief Provides a list of all fixed-point element types for use in GTest + * typed tests. + * + * Example: + * ``` + * // Invokes all typed fixture tests for all fixed-width types in libcudf + * TYPED_TEST_CASE(MyTypedFixture, cudf::test::FixedPointTypes); + * ``` + **/ +using FixedPointTypes = cudf::test::Types; + /** * @brief Provides a list of all fixed-width element types for use in GTest * typed tests. @@ -192,7 +205,21 @@ using ListTypes = cudf::test::Types; * TYPED_TEST_CASE(MyTypedFixture, cudf::test::FixedWidthTypes); * ``` **/ -using FixedWidthTypes = Concat; +using FixedWidthTypes = Concat; + +/** + * @brief Provides a list of all fixed-width element types except for the + * fixed-point types for use in GTest typed tests. Certain tests written for + * fixed-width types don't work for fixed-point as fixed-point types aren't + * constructible from other fixed-width types (a scale needs to be specified) + * + * Example: + * ``` + * // Invokes all typed fixture tests for all fixed-width types in libcudf + * TYPED_TEST_CASE(MyTypedFixture, cudf::test::FixedWidthTypes); + * ``` + **/ +using FixedWidthTypesWithoutFixedPoint = Concat; /** * @brief Provides a list of sortable types for use in GTest typed tests. diff --git a/cpp/tests/utilities_tests/column_utilities_tests.cpp b/cpp/tests/utilities_tests/column_utilities_tests.cpp index dea50974374..1cc8174ea58 100644 --- a/cpp/tests/utilities_tests/column_utilities_tests.cpp +++ b/cpp/tests/utilities_tests/column_utilities_tests.cpp @@ -44,7 +44,7 @@ template struct ColumnUtilitiesTestFloatingPoint : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(ColumnUtilitiesTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(ColumnUtilitiesTest, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST_CASE(ColumnUtilitiesTestIntegral, cudf::test::IntegralTypes); TYPED_TEST_CASE(ColumnUtilitiesTestFloatingPoint, cudf::test::FloatingPointTypes); diff --git a/cpp/tests/utilities_tests/column_wrapper_tests.cpp b/cpp/tests/utilities_tests/column_wrapper_tests.cpp index f5e2b2496c3..0cef36a88a4 100644 --- a/cpp/tests/utilities_tests/column_wrapper_tests.cpp +++ b/cpp/tests/utilities_tests/column_wrapper_tests.cpp @@ -30,7 +30,7 @@ struct FixedWidthColumnWrapperTest : public cudf::test::BaseFixture, auto data_type() { return cudf::data_type{cudf::type_to_id()}; } }; -TYPED_TEST_CASE(FixedWidthColumnWrapperTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(FixedWidthColumnWrapperTest, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(FixedWidthColumnWrapperTest, EmptyIterator) {