From eca134904dcf336143127a85123e0fde0a5a8250 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Fri, 12 Jun 2020 14:13:19 +0000 Subject: [PATCH 01/64] Initial changes --- cpp/include/cudf/detail/reduction.cuh | 5 ++- cpp/include/cudf/fixed_point/fixed_point.hpp | 2 ++ cpp/include/cudf/types.hpp | 1 + cpp/include/cudf/utilities/traits.hpp | 35 ++++++++++++++++++- .../cudf/utilities/type_dispatcher.hpp | 11 ++++++ 5 files changed, 50 insertions(+), 4 deletions(-) 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/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index 7063cdeec2d..b674994b3d5 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -771,5 +771,7 @@ std::ostream& operator<<(std::ostream& os, fixed_point const& fp) return os << static_cast(fp); } +using decimal32 = fixed_point; + /** @} */ // end of group } // namespace numeric diff --git a/cpp/include/cudf/types.hpp b/cpp/include/cudf/types.hpp index 5eb4234bec7..7a3420726f5 100644 --- a/cpp/include/cudf/types.hpp +++ b/cpp/include/cudf/types.hpp @@ -207,6 +207,7 @@ enum type_id { DICTIONARY32, ///< Dictionary type using int32 indices STRING, ///< String elements LIST, ///< List elements + DECIMAL32, ///< Fixed-point type with int32_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 4faa861a803..6ac912860c5 100644 --- a/cpp/include/cudf/utilities/traits.hpp +++ b/cpp/include/cudf/utilities/traits.hpp @@ -253,6 +253,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. * @@ -315,7 +348,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 ceb95051116..7303a839598 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,7 @@ 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); template struct type_to_scalar_type_impl { @@ -171,6 +174,12 @@ struct type_to_scalar_type_impl { using ScalarDeviceType = cudf::string_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; + using ScalarDeviceType = cudf::numeric_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; @@ -387,6 +396,8 @@ CUDA_HOST_DEVICE_CALLABLE constexpr decltype(auto) type_dispatcher(cudf::data_ty std::forward(args)...); case LIST: return f.template operator()::type>(std::forward(args)...); + case DECIMAL32: + return f.template operator()::type>(std::forward(args)...); default: { #ifndef __CUDA_ARCH__ CUDF_FAIL("Unsupported type_id."); From 6b2a3d185cd6a4035a6d97936e5740b086c8ed47 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 23 Jun 2020 23:38:12 +0000 Subject: [PATCH 02/64] Changes to support fixed_point_scalar --- cpp/include/cudf/scalar/scalar.hpp | 50 +++++++++++++++++++ .../cudf/scalar/scalar_device_view.cuh | 12 +++++ cpp/include/cudf/types.hpp | 8 +++ .../cudf/utilities/type_dispatcher.hpp | 4 +- 4 files changed, 72 insertions(+), 2 deletions(-) diff --git a/cpp/include/cudf/scalar/scalar.hpp b/cpp/include/cudf/scalar/scalar.hpp index 684f480b3b3..fd67035a56b 100644 --- a/cpp/include/cudf/scalar/scalar.hpp +++ b/cpp/include/cudf/scalar/scalar.hpp @@ -245,6 +245,56 @@ 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 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 + */ + 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 + */ + 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 * diff --git a/cpp/include/cudf/scalar/scalar_device_view.cuh b/cpp/include/cudf/scalar/scalar_device_view.cuh index 88328cb267c..ac006ad589d 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 7a3420726f5..ab9673eebe4 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; diff --git a/cpp/include/cudf/utilities/type_dispatcher.hpp b/cpp/include/cudf/utilities/type_dispatcher.hpp index 7303a839598..2e78aaf81da 100644 --- a/cpp/include/cudf/utilities/type_dispatcher.hpp +++ b/cpp/include/cudf/utilities/type_dispatcher.hpp @@ -176,8 +176,8 @@ struct type_to_scalar_type_impl { template <> // TODO: this is a temporary solution for make_pair_iterator struct type_to_scalar_type_impl { - using ScalarType = cudf::numeric_scalar; - using ScalarDeviceType = cudf::numeric_scalar_device_view; + 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 From 46de596f48876453dcace3fe5f17d4b78351f4c4 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 23 Jun 2020 23:38:39 +0000 Subject: [PATCH 03/64] Use default ctor in `replace` --- cpp/src/replace/replace.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/replace/replace.cu b/cpp/src/replace/replace.cu index f07c8cb118d..22d295f7e6d 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) { From 55645b8c03458e2aac9d74d558e58d6c8ab9aff3 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 24 Jun 2020 19:24:52 +0000 Subject: [PATCH 04/64] Add #pragma once --- cpp/include/cudf/fixed_point/fixed_point.hpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index b674994b3d5..9ffa3692495 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -14,6 +14,8 @@ * limitations under the License. */ +#pragma once + #include #include #include From 6d0f39e38a724d2355af6d54c215265040f0e578 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 24 Jun 2020 19:26:16 +0000 Subject: [PATCH 05/64] Use default ctor instead of 0 --- cpp/src/io/utilities/parsing_utils.cuh | 2 +- cpp/src/scalar/scalar_factories.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/io/utilities/parsing_utils.cuh b/cpp/src/io/utilities/parsing_utils.cuh index 90cf9ac2047..d69353b662b 100644 --- a/cpp/src/io/utilities/parsing_utils.cuh +++ b/cpp/src/io/utilities/parsing_utils.cuh @@ -142,7 +142,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/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); } From 3b1c4d25fd7bc251736f3777166adaba8b595cb0 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 24 Jun 2020 19:27:22 +0000 Subject: [PATCH 06/64] json/csv decode --- cpp/src/io/csv/csv_gpu.cu | 12 ++++++++++++ cpp/src/io/json/json_gpu.cu | 12 ++++++++++++ 2 files changed, 24 insertions(+) diff --git a/cpp/src/io/csv/csv_gpu.cu b/cpp/src/io/csv/csv_gpu.cu index 9e91f1dff48..6bd53da40bd 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,17 @@ __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{}; +} + /** * @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 e5b915112db..1b96e85c8ee 100644 --- a/cpp/src/io/json/json_gpu.cu +++ b/cpp/src/io/json/json_gpu.cu @@ -25,6 +25,7 @@ #include #include +#include #include #include @@ -253,6 +254,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, long start, @@ -261,6 +263,7 @@ __inline__ __device__ cudf::dictionary32 decode_value(const char *data, { return cudf::dictionary32{}; } + template <> __inline__ __device__ cudf::list_view decode_value(const char *data, long start, @@ -270,6 +273,15 @@ __inline__ __device__ cudf::list_view decode_value(const char *data, return cudf::list_view{}; } +template <> +__inline__ __device__ numeric::decimal32 decode_value(const char *data, + long start, + long end, + ParseOptions const &opts) +{ + return numeric::decimal32{}; +} + /** * @brief Functor for converting plain text data to cuDF data type value. **/ From 3045ddb90caf5ae2b97b0809988cb5b483bed979 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 24 Jun 2020 19:28:16 +0000 Subject: [PATCH 07/64] column_view_printer temoprary () operator --- cpp/tests/utilities/column_utilities.cu | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index 021c13a6ddb..9f559b22d1a 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -427,6 +427,19 @@ struct column_view_printer { this->template operator()(*col_as_strings, out, indent); } + template ::value>* = nullptr> + void operator()(cudf::column_view const& col, + std::vector& out, + std::string const& indent) + { + out.resize(col.size()); + std::transform(thrust::make_counting_iterator(size_type{0}), + thrust::make_counting_iterator(col.size()), + out.begin(), + [](auto idx) { return std::string(":p"); }); // #TODO + } + template ::value>* = nullptr> void operator()(cudf::column_view const& col, From a056bc516d01dd5c4c20797bd22d9eeb4e036c5a Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 24 Jun 2020 19:28:47 +0000 Subject: [PATCH 08/64] Changes for algorithms to compile --- .../cudf/detail/aggregation/aggregation.cuh | 35 ++++++++++--------- cpp/src/copying/copy.cu | 19 ++++++++++ cpp/src/copying/get_element.cu | 12 ++++++- cpp/src/replace/clamp.cu | 13 +++++++ 4 files changed, 62 insertions(+), 17 deletions(-) diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index 0be3400faf8..2270cb82494 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() and not cudf::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/src/copying/copy.cu b/cpp/src/copying/copy.cu index ef05605ad4e..8c30daa431b 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,24 @@ struct copy_if_else_functor_impl { } }; +/** + * @brief Specialization of copy_if_else_functor for fixed_point. + */ +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 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..966e71902c9 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,16 @@ 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"); + } }; } // namespace diff --git a/cpp/src/replace/clamp.cu b/cpp/src/replace/clamp.cu index d9b2c50c081..a53e7310703 100644 --- a/cpp/src/replace/clamp.cu +++ b/cpp/src/replace/clamp.cu @@ -286,6 +286,19 @@ 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"); +} + /** * @copydoc cudf::clamp(column_view const& input, scalar const& lo, From e05591e00ee5f68475d54e33c0e71bdf72be248c Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 24 Jun 2020 21:49:49 +0000 Subject: [PATCH 09/64] Add TOREVIEW comment --- cpp/include/cudf/detail/utilities/device_operators.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cudf/detail/utilities/device_operators.cuh b/cpp/include/cudf/detail/utilities/device_operators.cuh index e2414e6c5d5..b76e8ba01d9 100644 --- a/cpp/include/cudf/detail/utilities/device_operators.cuh +++ b/cpp/include/cudf/detail/utilities/device_operators.cuh @@ -97,7 +97,7 @@ struct DeviceMin { typename std::enable_if_t::value>* = nullptr> static constexpr T identity() { - return std::numeric_limits::max(); + return std::numeric_limits::max(); // #TODO #TOREVIEW } // @brief identity specialized for string_view From 7c3efda58d9cee8e6790481b43fc61ceb88a48a9 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 24 Jun 2020 21:50:07 +0000 Subject: [PATCH 10/64] Add sorted_order and gather tests --- cpp/tests/fixed_point/fixed_point_tests.cu | 44 ++++++++++++++++++++-- 1 file changed, 41 insertions(+), 3 deletions(-) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 29cf87a8e13..4f89c84db9b 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -14,14 +14,23 @@ * limitations under the License. */ -#include +#include +#include +#include +#include "tests/utilities/column_utilities.hpp" +#include "tests/utilities/table_utilities.hpp" + #include +#include #include +#include +#include #include +#include "cudf/types.hpp" + +#include #include #include -#include -#include #include #include @@ -514,4 +523,33 @@ TEST_F(FixedPointTest, DecimalXXThrustOnDevice) EXPECT_EQ(vec2, vec3); } +template +using wrapper = cudf::test::fixed_width_column_wrapper; + +TEST_F(FixedPointTest, FixedPointSortedOrderGather) +{ + auto const ZERO = decimal32{0, scale_type{0}}; + auto const ONE = decimal32{1, scale_type{0}}; + auto const TWO = decimal32{2, scale_type{0}}; + auto const THREE = decimal32{3, scale_type{0}}; + auto const FOUR = decimal32{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()); +} + CUDF_TEST_PROGRAM_MAIN() From 824fb37c8cb3caa943f56dd4266518d4b28f9cf4 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 24 Jun 2020 21:50:35 +0000 Subject: [PATCH 11/64] Remove column_device_printer hack --- cpp/tests/utilities/column_utilities.cu | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index 9f559b22d1a..afddddc06e8 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -433,11 +433,16 @@ struct column_view_printer { 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) { return std::string(":p"); }); // #TODO + [&](auto idx) { + auto const d = static_cast(h_data.first[idx]); + return std::to_string(d); + }); } template Date: Mon, 29 Jun 2020 14:45:41 +0000 Subject: [PATCH 12/64] Small cleanups --- cpp/src/binaryop/binaryop.cpp | 17 +++++++++-------- cpp/src/jit/type.cpp | 3 +-- 2 files changed, 10 insertions(+), 10 deletions(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index be368a57088..99822f3a07e 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -325,14 +325,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/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) From e0f51f842114e288f271263a27f59c70a0d93940 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 9 Jul 2020 01:01:55 +0000 Subject: [PATCH 13/64] Remove unused header --- cpp/include/cudf/fixed_point/fixed_point.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index 9ffa3692495..eec8d4ee655 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -18,7 +18,6 @@ #include #include -#include #include #include From 29101abfa549e094f6a4c436c0dfda0413421826 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 9 Jul 2020 02:54:05 +0000 Subject: [PATCH 14/64] Add missing header --- cpp/include/cudf/fixed_point/fixed_point.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index eec8d4ee655..7a05f6fa9f0 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -20,6 +20,7 @@ #include #include #include +#include //! `fixed_point` and supporting types namespace numeric { From d83fcb2ed010577abb812301b9a770d531e9f468 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 9 Jul 2020 16:16:16 +0000 Subject: [PATCH 15/64] Remove << ostream operator / iostream --- cpp/include/cudf/fixed_point/fixed_point.hpp | 16 +--------------- 1 file changed, 1 insertion(+), 15 deletions(-) diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index 7a05f6fa9f0..b77c1d94658 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -18,8 +18,8 @@ #include #include -#include #include +#include #include //! `fixed_point` and supporting types @@ -759,20 +759,6 @@ 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; /** @} */ // end of group From 8fe05c6ad9b41d1c5f33367471cabe58fc6073ed Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 9 Jul 2020 16:16:56 +0000 Subject: [PATCH 16/64] Add changes to support binaryops + unit test (failing) --- cpp/CMakeLists.txt | 23 ++++++++++++++++- cpp/include/cudf/column/column_factories.hpp | 2 +- cpp/src/binaryop/binaryop.cpp | 12 ++++++--- cpp/src/binaryop/jit/code/kernel.cpp | 1 + cpp/src/jit/common_headers.hpp | 6 +++++ cpp/tests/fixed_point/fixed_point_tests.cu | 27 ++++++++++++++++++++ 6 files changed, 66 insertions(+), 5 deletions(-) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 3f14909ad38..b70f1b5e4f0 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -617,6 +617,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/libcudacxx/details/__config.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/limits.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cfloat.jit @@ -625,6 +626,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 @@ -633,10 +636,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 ${LIBCUDACXX_INCLUDE_DIR}/details/__config ${LIBCUDACXX_INCLUDE_DIR}/simt/limits ${LIBCUDACXX_INCLUDE_DIR}/simt/cfloat @@ -645,6 +651,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 @@ -653,6 +661,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 @@ -660,6 +670,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 ${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 COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/cfloat libcudacxx_simt_cfloat > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cfloat.jit @@ -667,6 +678,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 @@ -676,6 +689,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 @@ -683,6 +698,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/libcudacxx/details/__config.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/limits.jit ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cfloat.jit @@ -690,6 +706,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 @@ -698,7 +716,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/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 99822f3a07e..102acbed3a4 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -36,11 +36,13 @@ #include "compiled/binary_ops.hpp" #include +#include #include -#include #include #include +#include + namespace cudf { namespace binops { @@ -69,8 +71,12 @@ namespace jit { const std::string hash = "prog_binop"; -const std::vector header_names{ - "operation.h", "traits.h", cudf_types_hpp, cudf_utilities_bit_hpp, cudf_wrappers_timestamps_hpp}; +const std::vector header_names{"operation.h", + "traits.h", + cudf_types_hpp, + cudf_utilities_bit_hpp, + cudf_wrappers_timestamps_hpp, + cudf_fixed_point_fixed_point_hpp}; std::istream* headers_code(std::string filename, std::iostream& stream) { diff --git a/cpp/src/binaryop/jit/code/kernel.cpp b/cpp/src/binaryop/jit/code/kernel.cpp index 2ca2049f4db..a6217640b04 100644 --- a/cpp/src/binaryop/jit/code/kernel.cpp +++ b/cpp/src/binaryop/jit/code/kernel.cpp @@ -30,6 +30,7 @@ const char* kernel = #include #include #include + #include #include "operation.h" template diff --git a/cpp/src/jit/common_headers.hpp b/cpp/src/jit/common_headers.hpp index 4ea7867fc7f..4c680e0e64a 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 @@ -67,6 +71,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/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 4f89c84db9b..c04d0a6f095 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -20,6 +20,7 @@ #include "tests/utilities/column_utilities.hpp" #include "tests/utilities/table_utilities.hpp" +#include #include #include #include @@ -552,4 +553,30 @@ TEST_F(FixedPointTest, FixedPointSortedOrderGather) cudf::test::expect_tables_equal(sorted_table, sorted->view()); } +TEST_F(FixedPointTest, FixedPointBinaryOpAdd) +{ + auto const sz = std::size_t{1000}; + + auto vec1 = std::vector(sz); + auto const vec2 = std::vector(sz, decimal32{1, scale_type{-1}}); + auto expected = std::vector(sz); + + std::iota(std::begin(vec1), std::end(vec1), decimal32{}); + + 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, cudf::data_type(cudf::type_id::DECIMAL32)); + + cudf::test::expect_columns_equal(expected_col, result->view()); +} + CUDF_TEST_PROGRAM_MAIN() From f2b0b8ab30e553aa3dbaa3f89b99ae57a0c53da5 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Fri, 10 Jul 2020 17:57:57 +0000 Subject: [PATCH 17/64] Final changes for `binaryop` to work --- cpp/include/cudf/fixed_point/fixed_point.hpp | 47 +++++++++++--------- cpp/include/cudf/utilities/traits.hpp | 2 +- 2 files changed, 27 insertions(+), 22 deletions(-) diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index b77c1d94658..7ead29791a1 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -16,11 +16,16 @@ #pragma once +#define _LIBCUDACXX_USE_CXX17_TYPE_TRAITS + +#include +#include // add simt namespace + +#include #include #include -#include + #include -#include //! `fixed_point` and supporting types namespace numeric { @@ -53,13 +58,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 @@ -80,8 +85,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); @@ -202,7 +207,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 { @@ -230,7 +235,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 { @@ -258,7 +263,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; @@ -288,8 +293,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} { @@ -318,7 +323,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)); @@ -564,9 +569,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"; @@ -583,8 +588,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 @@ -598,8 +603,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 @@ -613,7 +618,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 @@ -627,8 +632,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) diff --git a/cpp/include/cudf/utilities/traits.hpp b/cpp/include/cudf/utilities/traits.hpp index 15b9c396592..fc07b38a244 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 { From ebb929fee8384619949db6603fb386c6bd541483 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 13 Jul 2020 14:13:31 -0400 Subject: [PATCH 18/64] Add unit test for fixed_point concatenate --- cpp/tests/fixed_point/fixed_point_tests.cu | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index c04d0a6f095..3e9ef53b774 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -579,4 +579,20 @@ TEST_F(FixedPointTest, FixedPointBinaryOpAdd) cudf::test::expect_columns_equal(expected_col, result->view()); } +TEST_F(FixedPointTest, FixedPointConcatentate) +{ + auto vec = std::vector(1000); + std::iota(std::begin(vec), std::end(vec), decimal32{}); + + 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); +} + CUDF_TEST_PROGRAM_MAIN() From 9544fa02ef584f095f52a461b12c71d562d1ad9d Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 13 Jul 2020 23:05:39 -0400 Subject: [PATCH 19/64] Add unit test for fixed_point replace --- cpp/tests/fixed_point/fixed_point_tests.cu | 25 ++++++++++++++++++++++ 1 file changed, 25 insertions(+) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 3e9ef53b774..3e0b98026b7 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -24,6 +24,7 @@ #include #include #include +#include #include #include #include @@ -595,4 +596,28 @@ TEST_F(FixedPointTest, FixedPointConcatentate) cudf::test::expect_columns_equal(*results, expected); } +TEST_F(FixedPointTest, FixedPointReplace) +{ + auto const ONE = decimal32{1, scale_type{0}}; + auto const TWO = decimal32{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, mr()); + + cudf::test::expect_columns_equal(*result, expected_w); +} + CUDF_TEST_PROGRAM_MAIN() From 1c4a00772c824a63fed317db5440ef358a215edd Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 14 Jul 2020 10:32:51 -0400 Subject: [PATCH 20/64] Add unit tests for fixed_point upper_bound & lower_bound --- cpp/tests/fixed_point/fixed_point_tests.cu | 41 ++++++++++++++++++++++ 1 file changed, 41 insertions(+) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 3e0b98026b7..bc4f6f656de 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -28,6 +28,7 @@ #include #include #include +#include "cudf/search.hpp" #include "cudf/types.hpp" #include @@ -620,4 +621,44 @@ TEST_F(FixedPointTest, FixedPointReplace) cudf::test::expect_columns_equal(*result, expected_w); } +TEST_F(FixedPointTest, FixedPointLowerBound) +{ + auto vec = std::vector(1000); + std::iota(std::begin(vec), std::end(vec), decimal32{}); + + auto const values = wrapper{decimal32{200, scale_type{0}}, + decimal32{400, scale_type{0}}, + decimal32{600, scale_type{0}}, + decimal32{800, scale_type{0}}}; + auto const expect = wrapper{200, 400, 600, 800}; + auto const 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); +} + +TEST_F(FixedPointTest, FixedPointUpperBound) +{ + auto vec = std::vector(1000); + std::iota(std::begin(vec), std::end(vec), decimal32{}); + + auto const values = wrapper{decimal32{200, scale_type{0}}, + decimal32{400, scale_type{0}}, + decimal32{600, scale_type{0}}, + decimal32{800, scale_type{0}}}; + auto const expect = wrapper{201, 401, 601, 801}; + auto const 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() From e69c2c7cf058b7c5d894da8d60b18e78d7d84c03 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 15 Jul 2020 09:36:24 -0400 Subject: [PATCH 21/64] Remove trailing comma --- cpp/tests/reshape/interleave_columns_tests.cpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index ab3beea9fc1..54b2e50af33 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -60,10 +60,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); From ec9fbd6e27318f85129d2d0aceb9ab8fa55fba98 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 15 Jul 2020 09:41:53 -0400 Subject: [PATCH 22/64] Add unit test for fixed_point interleave --- cpp/tests/fixed_point/fixed_point_tests.cu | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index bc4f6f656de..2bf740b3337 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -25,6 +25,7 @@ #include #include #include +#include #include #include #include @@ -661,4 +662,21 @@ TEST_F(FixedPointTest, FixedPointUpperBound) expect_columns_equal(*result, expect); } +TEST_F(FixedPointTest, FixedPointInterleave) +{ + auto const ONE = decimal32{1, scale_type{0}}; + auto const TWO = decimal32{2, scale_type{0}}; + auto const THREE = decimal32{3, scale_type{0}}; + auto const FOUR = decimal32{4, scale_type{0}}; + + auto const a = wrapper({ONE, THREE}); + auto const b = wrapper({TWO, FOUR}); + + auto const input = cudf::table_view{std::vector{a, b}}; + auto const expected = wrapper({ONE, TWO, THREE, FOUR}); + auto const actual = cudf::interleave_columns(input); + + expect_columns_equal(expected, actual->view()); +} + CUDF_TEST_PROGRAM_MAIN() From b2ef6039cd703523911fede603eedbc330d0f7ff Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 15 Jul 2020 09:42:53 -0400 Subject: [PATCH 23/64] Add "scale loop" to interleave unit test --- cpp/tests/fixed_point/fixed_point_tests.cu | 22 ++++++++++++---------- 1 file changed, 12 insertions(+), 10 deletions(-) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 2bf740b3337..08305f67f3f 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -664,19 +664,21 @@ TEST_F(FixedPointTest, FixedPointUpperBound) TEST_F(FixedPointTest, FixedPointInterleave) { - auto const ONE = decimal32{1, scale_type{0}}; - auto const TWO = decimal32{2, scale_type{0}}; - auto const THREE = decimal32{3, scale_type{0}}; - auto const FOUR = decimal32{4, scale_type{0}}; + for (int i = 0; i > -4; --i) { + auto const ONE = decimal32{1, scale_type{i}}; + auto const TWO = decimal32{2, scale_type{i}}; + auto const THREE = decimal32{3, scale_type{i}}; + auto const FOUR = decimal32{4, scale_type{i}}; - auto const a = wrapper({ONE, THREE}); - auto const b = wrapper({TWO, FOUR}); + auto const a = wrapper({ONE, THREE}); + auto const b = wrapper({TWO, FOUR}); - auto const input = cudf::table_view{std::vector{a, b}}; - auto const expected = wrapper({ONE, TWO, THREE, FOUR}); - auto const actual = cudf::interleave_columns(input); + auto const input = cudf::table_view{std::vector{a, b}}; + auto const expected = wrapper({ONE, TWO, THREE, FOUR}); + auto const actual = cudf::interleave_columns(input); - expect_columns_equal(expected, actual->view()); + expect_columns_equal(expected, actual->view()); + } } CUDF_TEST_PROGRAM_MAIN() From 5ca44bdc25b812580a7ea6e1ecfe2d0897060cd6 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 15 Jul 2020 16:55:44 -0400 Subject: [PATCH 24/64] Add unit test for fixed_point product reduction --- .../detail/utilities/device_operators.cuh | 14 +++++++++++-- cpp/src/reductions/simple.cuh | 4 +++- cpp/tests/fixed_point/fixed_point_tests.cu | 21 +++++++++++++++++++ 3 files changed, 36 insertions(+), 3 deletions(-) diff --git a/cpp/include/cudf/detail/utilities/device_operators.cuh b/cpp/include/cudf/detail/utilities/device_operators.cuh index b76e8ba01d9..150310922ec 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 { @@ -49,7 +53,7 @@ struct DeviceSum { template static constexpr T identity() { - return T{0}; + return T{}; } }; @@ -157,11 +161,17 @@ struct DeviceProduct { return lhs * rhs; } - template + template ()>* = nullptr> static constexpr T identity() { return T{1}; } + + template ()>* = nullptr> + static constexpr T identity() + { + return T{1, numeric::scale_type{0}}; + } }; /* @brief binary `and` operator */ diff --git a/cpp/src/reductions/simple.cuh b/cpp/src/reductions/simple.cuh index 0aff7328fca..f8ea60c3f64 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 std::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/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 08305f67f3f..de6cba6fe42 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -24,6 +24,7 @@ #include #include #include +#include #include #include #include @@ -681,4 +682,24 @@ TEST_F(FixedPointTest, FixedPointInterleave) } } +TEST_F(FixedPointTest, FixedPointReductionProduct) +{ + auto const ONE = decimal32{1, scale_type{0}}; + auto const TWO = decimal32{2, scale_type{0}}; + auto const THREE = decimal32{3, scale_type{0}}; + auto const FOUR = decimal32{4, scale_type{0}}; + auto const _24 = decimal32{24, scale_type{0}}; + + auto const in = std::vector{ONE, TWO, THREE, FOUR}; + auto const column = wrapper(in.cbegin(), in.cend()); + auto const expected = std::accumulate(in.cbegin(), in.cend(), ONE, std::multiplies()); + auto const out_type = cudf::data_type{cudf::type_id::DECIMAL32}; + + 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); +} + CUDF_TEST_PROGRAM_MAIN() From 1a6f1317395b1290d8158ec8167e3d4b8bffd775 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 15 Jul 2020 16:59:47 -0400 Subject: [PATCH 25/64] Add unit test for fixed_point sum reduction --- cpp/tests/fixed_point/fixed_point_tests.cu | 21 +++++++++++++++++++++ 1 file changed, 21 insertions(+) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index de6cba6fe42..caa2d7a4f1a 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -702,4 +702,25 @@ TEST_F(FixedPointTest, FixedPointReductionProduct) EXPECT_EQ(result_scalar->value(), _24); } +TEST_F(FixedPointTest, FixedPointReductionSum) +{ + auto const ZERO = decimal32{0, scale_type{0}}; + auto const ONE = decimal32{1, scale_type{0}}; + auto const TWO = decimal32{2, scale_type{0}}; + auto const THREE = decimal32{3, scale_type{0}}; + auto const FOUR = decimal32{4, scale_type{0}}; + auto const TEN = decimal32{10, scale_type{0}}; + + auto const in = std::vector{ONE, TWO, THREE, FOUR}; + auto const column = wrapper(in.cbegin(), in.cend()); + auto const expected = std::accumulate(in.cbegin(), in.cend(), ZERO, std::plus()); + auto const out_type = cudf::data_type{cudf::type_id::DECIMAL32}; + + 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() From c54488df7b9a14b2aaf1426c484d5e8ab3d1af59 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 15 Jul 2020 21:03:06 -0400 Subject: [PATCH 26/64] Update CHANGELOG --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 23ad5d675f9..364e5889331 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -42,6 +42,7 @@ - PR #5623 Add `is_ipv4` strings 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 `fixed_point` Column Support (Part 1) ## Improvements From a99e34eace92ada5df21058ba177fc93c4d7b14d Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 16 Jul 2020 11:48:01 -0400 Subject: [PATCH 27/64] Add decimal64 alias --- cpp/include/cudf/fixed_point/fixed_point.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index 7ead29791a1..f77d3fb06ae 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -765,6 +765,7 @@ CUDA_HOST_DEVICE_CALLABLE bool operator>(fixed_point const& lhs, } using decimal32 = fixed_point; +using decimal64 = fixed_point; /** @} */ // end of group } // namespace numeric From 661a5ed1bf41372905bf02753dc70ad153f98dd6 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 16 Jul 2020 12:14:07 -0400 Subject: [PATCH 28/64] Add decimal64 to fixed_point type trait --- cpp/include/cudf/utilities/traits.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cudf/utilities/traits.hpp b/cpp/include/cudf/utilities/traits.hpp index fc07b38a244..66907df62bb 100644 --- a/cpp/include/cudf/utilities/traits.hpp +++ b/cpp/include/cudf/utilities/traits.hpp @@ -339,7 +339,7 @@ constexpr inline bool is_timestamp(data_type type) template constexpr inline bool is_fixed_point() { - return std::is_same::value; // || std::is_same::value; + return std::is_same::value || std::is_same::value; } struct is_fixed_point_impl { From 4e3d35b9b8f405213cfb8ae17bb4911bb7570dbd Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 16 Jul 2020 12:20:11 -0400 Subject: [PATCH 29/64] Add decimal64 to types.hpp and type_dispatcher.hpp --- cpp/include/cudf/types.hpp | 1 + cpp/include/cudf/utilities/type_dispatcher.hpp | 12 +++++++++++- 2 files changed, 12 insertions(+), 1 deletion(-) diff --git a/cpp/include/cudf/types.hpp b/cpp/include/cudf/types.hpp index 3745d262a41..f19ce096efb 100644 --- a/cpp/include/cudf/types.hpp +++ b/cpp/include/cudf/types.hpp @@ -216,6 +216,7 @@ enum class type_id : int32_t { 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/type_dispatcher.hpp b/cpp/include/cudf/utilities/type_dispatcher.hpp index 7ddf1c55b89..0f9f61a91f5 100644 --- a/cpp/include/cudf/utilities/type_dispatcher.hpp +++ b/cpp/include/cudf/utilities/type_dispatcher.hpp @@ -135,6 +135,7 @@ 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 { @@ -174,12 +175,18 @@ struct type_to_scalar_type_impl { using ScalarDeviceType = cudf::string_scalar_device_view; }; -template <> // TODO: this is a temporary solution for make_pair_iterator +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; @@ -412,6 +419,9 @@ CUDA_HOST_DEVICE_CALLABLE constexpr decltype(auto) type_dispatcher(cudf::data_ty 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."); From 51b8c0240bc1cc3748e1e4ea58f0e542d0447f59 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 16 Jul 2020 12:44:55 -0400 Subject: [PATCH 30/64] Make column_device_printer generic for decimal64 --- cpp/tests/utilities/column_utilities.cu | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index a2d1db54a72..3e25f7ab08b 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -427,8 +427,7 @@ struct column_view_printer { this->template operator()(*col_as_strings, out, indent); } - template ::value>* = nullptr> + template ()>* = nullptr> void operator()(cudf::column_view const& col, std::vector& out, std::string const& indent) From 3b30ff7a34e1558f768de5a6d133a62f99f6b441 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 16 Jul 2020 12:46:05 -0400 Subject: [PATCH 31/64] Add code for stamping out decimal64 algorithm specializations --- cpp/src/copying/copy.cu | 20 +++++++++++++++++++- cpp/src/copying/get_element.cu | 10 ++++++++++ cpp/src/io/csv/csv_gpu.cu | 11 +++++++++++ cpp/src/io/json/json_gpu.cu | 9 +++++++++ cpp/src/replace/clamp.cu | 13 +++++++++++++ 5 files changed, 62 insertions(+), 1 deletion(-) diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu index 557ade901e9..2def3471729 100644 --- a/cpp/src/copying/copy.cu +++ b/cpp/src/copying/copy.cu @@ -117,7 +117,7 @@ struct copy_if_else_functor_impl { }; /** - * @brief Specialization of copy_if_else_functor for fixed_point. + * @brief Specialization of copy_if_else_functor for decimal32. */ template struct copy_if_else_functor_impl { @@ -134,6 +134,24 @@ struct copy_if_else_functor_impl { } }; +/** + * @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 966e71902c9..6408d1c7720 100644 --- a/cpp/src/copying/get_element.cu +++ b/cpp/src/copying/get_element.cu @@ -124,6 +124,16 @@ struct get_element_functor { { 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 6bd53da40bd..3ae1ef9405c 100644 --- a/cpp/src/io/csv/csv_gpu.cu +++ b/cpp/src/io/csv/csv_gpu.cu @@ -414,6 +414,17 @@ __inline__ __device__ numeric::decimal32 decode_value(const char *data, 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 2752ff6a3a8..28dc82accb5 100644 --- a/cpp/src/io/json/json_gpu.cu +++ b/cpp/src/io/json/json_gpu.cu @@ -282,6 +282,15 @@ __inline__ __device__ numeric::decimal32 decode_value(const char *data, return numeric::decimal32{}; } +template <> +__inline__ __device__ numeric::decimal64 decode_value(const char *data, + long start, + long end, + ParseOptions const &opts) +{ + return numeric::decimal64{}; +} + /** * @brief Functor for converting plain text data to cuDF data type value. **/ diff --git a/cpp/src/replace/clamp.cu b/cpp/src/replace/clamp.cu index a53e7310703..a47892baf39 100644 --- a/cpp/src/replace/clamp.cu +++ b/cpp/src/replace/clamp.cu @@ -299,6 +299,19 @@ std::unique_ptr dispatch_clamp::operator()( 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, From 51487d11a6ae998b60397fd3f319fc3f7c49653e Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 16 Jul 2020 13:19:18 -0400 Subject: [PATCH 32/64] Generalizing unit tests for decimal64 --- cpp/tests/fixed_point/fixed_point_tests.cu | 174 ++++++++++++--------- 1 file changed, 96 insertions(+), 78 deletions(-) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index caa2d7a4f1a..09594c8ab43 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -531,21 +531,23 @@ TEST_F(FixedPointTest, DecimalXXThrustOnDevice) template using wrapper = cudf::test::fixed_width_column_wrapper; -TEST_F(FixedPointTest, FixedPointSortedOrderGather) +TYPED_TEST(FixedPointTestBothReps, FixedPointSortedOrderGather) { - auto const ZERO = decimal32{0, scale_type{0}}; - auto const ONE = decimal32{1, scale_type{0}}; - auto const TWO = decimal32{2, scale_type{0}}; - auto const THREE = decimal32{3, scale_type{0}}; - auto const FOUR = decimal32{4, scale_type{0}}; + using decimalXX = fixed_point; + + 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 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 sorted_vec = std::vector{ZERO, ONE, TWO, THREE, FOUR}; - auto const input_col = wrapper(input_vec.begin(), input_vec.end()); + 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_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}}; @@ -559,6 +561,8 @@ TEST_F(FixedPointTest, FixedPointSortedOrderGather) TEST_F(FixedPointTest, FixedPointBinaryOpAdd) { + // using decimalXX = fixed_point; + auto const sz = std::size_t{1000}; auto vec1 = std::vector(sz); @@ -578,62 +582,68 @@ TEST_F(FixedPointTest, FixedPointBinaryOpAdd) auto const expected_col = wrapper(expected.begin(), expected.end()); auto const result = cudf::binary_operation( - lhs, rhs, cudf::binary_operator::ADD, cudf::data_type(cudf::type_id::DECIMAL32)); + lhs, rhs, cudf::binary_operator::ADD, static_cast(lhs).type()); cudf::test::expect_columns_equal(expected_col, result->view()); } -TEST_F(FixedPointTest, FixedPointConcatentate) +TYPED_TEST(FixedPointTestBothReps, FixedPointConcatentate) { - auto vec = std::vector(1000); - std::iota(std::begin(vec), std::end(vec), decimal32{}); + using decimalXX = fixed_point; + + 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 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()); + auto const expected = wrapper(vec.begin(), vec.end()); cudf::test::expect_columns_equal(*results, expected); } -TEST_F(FixedPointTest, FixedPointReplace) +TYPED_TEST(FixedPointTestBothReps, FixedPointReplace) { - auto const ONE = decimal32{1, scale_type{0}}; - auto const TWO = decimal32{2, scale_type{0}}; + using decimalXX = fixed_point; + + 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); + 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 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 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, mr()); + auto const result = cudf::find_and_replace_all(input_w, to_replace_w, replacement_w); cudf::test::expect_columns_equal(*result, expected_w); } -TEST_F(FixedPointTest, FixedPointLowerBound) +TYPED_TEST(FixedPointTestBothReps, FixedPointLowerBound) { - auto vec = std::vector(1000); - std::iota(std::begin(vec), std::end(vec), decimal32{}); + using decimalXX = fixed_point; + + auto vec = std::vector(1000); + std::iota(std::begin(vec), std::end(vec), decimalXX{}); - auto const values = wrapper{decimal32{200, scale_type{0}}, - decimal32{400, scale_type{0}}, - decimal32{600, scale_type{0}}, - decimal32{800, scale_type{0}}}; + auto const values = wrapper{decimalXX{200, scale_type{0}}, + decimalXX{400, scale_type{0}}, + decimalXX{600, scale_type{0}}, + decimalXX{800, scale_type{0}}}; auto const expect = wrapper{200, 400, 600, 800}; - auto const column = wrapper(vec.begin(), vec.end()); + auto const column = wrapper(vec.begin(), vec.end()); auto result = cudf::lower_bound({cudf::table_view{{column}}}, {cudf::table_view{{values}}}, @@ -643,17 +653,19 @@ TEST_F(FixedPointTest, FixedPointLowerBound) expect_columns_equal(*result, expect); } -TEST_F(FixedPointTest, FixedPointUpperBound) +TYPED_TEST(FixedPointTestBothReps, FixedPointUpperBound) { - auto vec = std::vector(1000); - std::iota(std::begin(vec), std::end(vec), decimal32{}); + using decimalXX = fixed_point; + + auto vec = std::vector(1000); + std::iota(std::begin(vec), std::end(vec), decimalXX{}); - auto const values = wrapper{decimal32{200, scale_type{0}}, - decimal32{400, scale_type{0}}, - decimal32{600, scale_type{0}}, - decimal32{800, scale_type{0}}}; + auto const values = wrapper{decimalXX{200, scale_type{0}}, + decimalXX{400, scale_type{0}}, + decimalXX{600, scale_type{0}}, + decimalXX{800, scale_type{0}}}; auto const expect = wrapper{201, 401, 601, 801}; - auto const column = wrapper(vec.begin(), vec.end()); + auto const column = wrapper(vec.begin(), vec.end()); auto result = cudf::upper_bound({cudf::table_view{{column}}}, {cudf::table_view{{values}}}, @@ -663,61 +675,67 @@ TEST_F(FixedPointTest, FixedPointUpperBound) expect_columns_equal(*result, expect); } -TEST_F(FixedPointTest, FixedPointInterleave) +TYPED_TEST(FixedPointTestBothReps, FixedPointInterleave) { + using decimalXX = fixed_point; + for (int i = 0; i > -4; --i) { - auto const ONE = decimal32{1, scale_type{i}}; - auto const TWO = decimal32{2, scale_type{i}}; - auto const THREE = decimal32{3, scale_type{i}}; - auto const FOUR = decimal32{4, scale_type{i}}; + auto const ONE = decimalXX{1, scale_type{i}}; + auto const TWO = decimalXX{2, scale_type{i}}; + auto const THREE = decimalXX{3, scale_type{i}}; + auto const FOUR = decimalXX{4, scale_type{i}}; - auto const a = wrapper({ONE, THREE}); - auto const b = wrapper({TWO, FOUR}); + auto const a = wrapper({ONE, THREE}); + auto const b = wrapper({TWO, FOUR}); auto const input = cudf::table_view{std::vector{a, b}}; - auto const expected = wrapper({ONE, TWO, THREE, FOUR}); + auto const expected = wrapper({ONE, TWO, THREE, FOUR}); auto const actual = cudf::interleave_columns(input); expect_columns_equal(expected, actual->view()); } } -TEST_F(FixedPointTest, FixedPointReductionProduct) +TYPED_TEST(FixedPointTestBothReps, FixedPointReductionProduct) { - auto const ONE = decimal32{1, scale_type{0}}; - auto const TWO = decimal32{2, scale_type{0}}; - auto const THREE = decimal32{3, scale_type{0}}; - auto const FOUR = decimal32{4, scale_type{0}}; - auto const _24 = decimal32{24, scale_type{0}}; + using decimalXX = fixed_point; - auto const in = std::vector{ONE, TWO, THREE, FOUR}; - auto const column = wrapper(in.cbegin(), in.cend()); - auto const expected = std::accumulate(in.cbegin(), in.cend(), ONE, std::multiplies()); - auto const out_type = cudf::data_type{cudf::type_id::DECIMAL32}; + 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 = 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(); auto const result = cudf::reduce(column, cudf::make_product_aggregation(), out_type); - auto const result_scalar = static_cast*>(result.get()); + auto const result_scalar = static_cast*>(result.get()); EXPECT_EQ(result_scalar->value(), expected); EXPECT_EQ(result_scalar->value(), _24); } -TEST_F(FixedPointTest, FixedPointReductionSum) +TYPED_TEST(FixedPointTestBothReps, FixedPointReductionSum) { - auto const ZERO = decimal32{0, scale_type{0}}; - auto const ONE = decimal32{1, scale_type{0}}; - auto const TWO = decimal32{2, scale_type{0}}; - auto const THREE = decimal32{3, scale_type{0}}; - auto const FOUR = decimal32{4, scale_type{0}}; - auto const TEN = decimal32{10, scale_type{0}}; - - auto const in = std::vector{ONE, TWO, THREE, FOUR}; - auto const column = wrapper(in.cbegin(), in.cend()); - auto const expected = std::accumulate(in.cbegin(), in.cend(), ZERO, std::plus()); - auto const out_type = cudf::data_type{cudf::type_id::DECIMAL32}; + using decimalXX = fixed_point; + + 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 = 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(); auto const result = cudf::reduce(column, cudf::make_sum_aggregation(), out_type); - auto const result_scalar = static_cast*>(result.get()); + auto const result_scalar = static_cast*>(result.get()); EXPECT_EQ(result_scalar->value(), expected); EXPECT_EQ(result_scalar->value(), TEN); From 58b76e2e3501842fbbad133ee5e591bcafaeb2ee Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 16 Jul 2020 17:44:08 -0400 Subject: [PATCH 33/64] Remove unnecessary #include fixed_point in binaryop kernel.cpp --- cpp/src/binaryop/jit/code/kernel.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/src/binaryop/jit/code/kernel.cpp b/cpp/src/binaryop/jit/code/kernel.cpp index a6217640b04..2ca2049f4db 100644 --- a/cpp/src/binaryop/jit/code/kernel.cpp +++ b/cpp/src/binaryop/jit/code/kernel.cpp @@ -30,7 +30,6 @@ const char* kernel = #include #include #include - #include #include "operation.h" template From 13b0dd28f276936b871091e88c4abc38f44c4862 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 16 Jul 2020 19:53:10 -0400 Subject: [PATCH 34/64] Remove unnecessary header_names entry for fixed_point --- cpp/src/binaryop/binaryop.cpp | 11 +++-------- 1 file changed, 3 insertions(+), 8 deletions(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 102acbed3a4..71c2055d14b 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -38,11 +38,10 @@ #include #include #include +#include #include #include -#include - namespace cudf { namespace binops { @@ -71,12 +70,8 @@ namespace jit { const std::string hash = "prog_binop"; -const std::vector header_names{"operation.h", - "traits.h", - cudf_types_hpp, - cudf_utilities_bit_hpp, - cudf_wrappers_timestamps_hpp, - cudf_fixed_point_fixed_point_hpp}; +const std::vector header_names{ + "operation.h", "traits.h", cudf_types_hpp, cudf_utilities_bit_hpp, cudf_wrappers_timestamps_hpp}; std::istream* headers_code(std::string filename, std::iostream& stream) { From 7759dfc3238c4eb9dcbaecf19b9301ed840a1e3a Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Fri, 17 Jul 2020 09:48:32 -0400 Subject: [PATCH 35/64] BROKEN: Add unit test for decimal64 binaryop --- cpp/tests/fixed_point/fixed_point_tests.cu | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 09594c8ab43..d02ca51b9b9 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -559,27 +559,27 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointSortedOrderGather) cudf::test::expect_tables_equal(sorted_table, sorted->view()); } -TEST_F(FixedPointTest, FixedPointBinaryOpAdd) +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd) { - // using decimalXX = fixed_point; + using decimalXX = fixed_point; auto const sz = std::size_t{1000}; - auto vec1 = std::vector(sz); - auto const vec2 = std::vector(sz, decimal32{1, scale_type{-1}}); - auto expected = std::vector(sz); + 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), decimal32{}); + 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()); + 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 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()); From 0256cd5b87439b873687b9e4a5b4ea736cc604a5 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Fri, 17 Jul 2020 13:17:58 -0400 Subject: [PATCH 36/64] Add back necessary jitify changes (seems to work w/o b/c of caching bug) --- cpp/src/binaryop/binaryop.cpp | 8 ++++++-- cpp/src/binaryop/jit/code/kernel.cpp | 1 + 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 71c2055d14b..f0d53adfedb 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -70,8 +70,12 @@ namespace jit { const std::string hash = "prog_binop"; -const std::vector header_names{ - "operation.h", "traits.h", cudf_types_hpp, cudf_utilities_bit_hpp, cudf_wrappers_timestamps_hpp}; +const std::vector header_names{"operation.h", + "traits.h", + cudf_types_hpp, + cudf_utilities_bit_hpp, + cudf_wrappers_timestamps_hpp, + cudf_fixed_point_fixed_point_hpp}; std::istream* headers_code(std::string filename, std::iostream& stream) { diff --git a/cpp/src/binaryop/jit/code/kernel.cpp b/cpp/src/binaryop/jit/code/kernel.cpp index 2ca2049f4db..a6217640b04 100644 --- a/cpp/src/binaryop/jit/code/kernel.cpp +++ b/cpp/src/binaryop/jit/code/kernel.cpp @@ -30,6 +30,7 @@ const char* kernel = #include #include #include + #include #include "operation.h" template From 4f5b459a830f87b6670003310d3388802b0fff5f Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Fri, 17 Jul 2020 13:18:44 -0400 Subject: [PATCH 37/64] Add binaryop MUL unit test for fixed_point --- cpp/tests/fixed_point/fixed_point_tests.cu | 28 ++++++++++++++++++++++ 1 file changed, 28 insertions(+) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index d02ca51b9b9..73fe8dfe83d 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -587,6 +587,34 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd) cudf::test::expect_columns_equal(expected_col, result->view()); } +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiply) +{ + using decimalXX = fixed_point; + + 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, FixedPointConcatentate) { using decimalXX = fixed_point; From ac4837f47ee3751c0fbe608be57f9d30316dc979 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 21 Jul 2020 18:14:51 -0400 Subject: [PATCH 38/64] Disable reduction algorithms / device identity functions --- .../detail/utilities/device_operators.cuh | 33 ++++++++++++++-- cpp/tests/fixed_point/fixed_point_tests.cu | 39 +++++++++++-------- 2 files changed, 51 insertions(+), 21 deletions(-) diff --git a/cpp/include/cudf/detail/utilities/device_operators.cuh b/cpp/include/cudf/detail/utilities/device_operators.cuh index 150310922ec..25478629da8 100644 --- a/cpp/include/cudf/detail/utilities/device_operators.cuh +++ b/cpp/include/cudf/detail/utilities/device_operators.cuh @@ -50,9 +50,16 @@ struct DeviceSum { return lhs + rhs; } - template + template ()>* = nullptr> + static constexpr T identity() + { + return T{}; + } + + template ()>* = nullptr> static constexpr T identity() { + CUDF_FAIL("fixed_point does not yet support device operator identity"); return T{}; } }; @@ -98,10 +105,18 @@ 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() { - return std::numeric_limits::max(); // #TODO #TOREVIEW + CUDF_FAIL("fixed_point does not yet support DeviceMin identity"); + return std::numeric_limits::max(); } // @brief identity specialized for string_view @@ -128,11 +143,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() @@ -170,6 +194,7 @@ struct DeviceProduct { template ()>* = nullptr> static constexpr T identity() { + CUDF_FAIL("fixed_point does not yet support DeviceProduct identity"); return T{1, numeric::scale_type{0}}; } }; diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 73fe8dfe83d..5ad6a0da076 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -732,41 +732,46 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointReductionProduct) 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 _24 = decimalXX{24, scale_type{0}}; - auto const in = std::vector{ONE, TWO, THREE, FOUR}; - auto const column = wrapper(in.cbegin(), in.cend()); - auto const expected = std::accumulate(in.cbegin(), in.cend(), ONE, std::multiplies()); + auto const in = std::vector{ONE, TWO, THREE, FOUR}; + auto const 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(); - auto const result = cudf::reduce(column, cudf::make_product_aggregation(), out_type); - auto const result_scalar = static_cast*>(result.get()); + EXPECT_THROW(cudf::reduce(column, cudf::make_product_aggregation(), out_type), cudf::logic_error); - EXPECT_EQ(result_scalar->value(), expected); - EXPECT_EQ(result_scalar->value(), _24); + // 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 decimalXX = fixed_point; - auto const ZERO = decimalXX{0, scale_type{0}}; + // 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 TEN = decimalXX{10, scale_type{0}}; - auto const in = std::vector{ONE, TWO, THREE, FOUR}; - auto const column = wrapper(in.cbegin(), in.cend()); - auto const expected = std::accumulate(in.cbegin(), in.cend(), ZERO, std::plus()); + auto const in = std::vector{ONE, TWO, THREE, FOUR}; + auto const 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(); - auto const result = cudf::reduce(column, cudf::make_sum_aggregation(), out_type); - auto const result_scalar = static_cast*>(result.get()); + 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); + // EXPECT_EQ(result_scalar->value(), expected); + // EXPECT_EQ(result_scalar->value(), TEN); } CUDF_TEST_PROGRAM_MAIN() From bfb68b4a219e596a72b5e9bdb717a7088d5c8470 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 21 Jul 2020 18:48:10 -0400 Subject: [PATCH 39/64] Add is_primitive_type trait --- .../cudf/detail/aggregation/aggregation.cuh | 35 +++++++++---------- cpp/include/cudf/utilities/traits.hpp | 14 ++++++++ cpp/src/copying/get_element.cu | 2 +- 3 files changed, 31 insertions(+), 20 deletions(-) diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index 2270cb82494..40c26d2656e 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.cuh +++ b/cpp/include/cudf/detail/aggregation/aggregation.cuh @@ -96,12 +96,11 @@ struct update_target_element { }; template -struct update_target_element< - Source, - aggregation::MIN, - target_has_nulls, - source_has_nulls, - std::enable_if_t() && !is_fixed_point()>> { +struct update_target_element()>> { __device__ void operator()(mutable_column_device_view target, size_type target_index, column_device_view source, @@ -118,12 +117,11 @@ struct update_target_element< }; template -struct update_target_element< - Source, - aggregation::MAX, - target_has_nulls, - source_has_nulls, - std::enable_if_t() && !is_fixed_point()>> { +struct update_target_element()>> { __device__ void operator()(mutable_column_device_view target, size_type target_index, column_device_view source, @@ -140,12 +138,11 @@ struct update_target_element< }; template -struct update_target_element< - Source, - aggregation::SUM, - target_has_nulls, - source_has_nulls, - std::enable_if_t() && !is_fixed_point()>> { +struct update_target_element()>> { __device__ void operator()(mutable_column_device_view target, size_type target_index, column_device_view source, @@ -362,7 +359,7 @@ struct identity_initializer { template static constexpr bool is_supported() { - return cudf::is_fixed_width() and not cudf::is_fixed_point() and + return cudf::is_primitive_type() 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/utilities/traits.hpp b/cpp/include/cudf/utilities/traits.hpp index 66907df62bb..4d74a0b4f55 100644 --- a/cpp/include/cudf/utilities/traits.hpp +++ b/cpp/include/cudf/utilities/traits.hpp @@ -449,6 +449,20 @@ constexpr inline bool is_fixed_width(data_type type) return cudf::type_dispatcher(type, is_fixed_width_impl{}); } +/** + * @brief Indicates whether the type `T` is a "primitive type". Currently this + * is fixed width types not including fixed point + * + * @tparam T The type to verify + * @return true `T` is a primitive type + * @return false `T` is not a primitive type + **/ +template +constexpr inline bool is_primitive_type() +{ + return cudf::is_fixed_width() && !cudf::is_fixed_point(); +} + /** * @brief Indicates whether the type `T` is a compound type. * diff --git a/cpp/src/copying/get_element.cu b/cpp/src/copying/get_element.cu index 6408d1c7720..d5746188a74 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 () && !is_fixed_point()> *p = nullptr> + template ()> *p = nullptr> std::unique_ptr operator()( column_view const &input, size_type index, From 58b09baadd40f61f7b252c281de266d8c5632337 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 22 Jul 2020 11:55:07 -0400 Subject: [PATCH 40/64] Fixing merge conflicts --- cpp/src/unary/cast_ops.cu | 16 +++++++++++++--- 1 file changed, 13 insertions(+), 3 deletions(-) diff --git a/cpp/src/unary/cast_ops.cu b/cpp/src/unary/cast_ops.cu index 3051e1f8980..e38316d01d2 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,7 @@ struct dispatch_unary_cast_from { dispatch_unary_cast_from(column_view inp) : input(inp) {} - template ()>* = nullptr> + template ()>* = nullptr> std::unique_ptr operator()(data_type type, rmm::mr::device_memory_resource* mr, cudaStream_t stream) @@ -159,6 +161,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, From 51c5e638d52838e07eaf281134be66245c9cc81f Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 22 Jul 2020 12:58:52 -0400 Subject: [PATCH 41/64] Remove unused type_traits header --- cpp/include/cudf/scalar/scalar.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/include/cudf/scalar/scalar.hpp b/cpp/include/cudf/scalar/scalar.hpp index 7e8f766cae2..59b2f60c975 100644 --- a/cpp/include/cudf/scalar/scalar.hpp +++ b/cpp/include/cudf/scalar/scalar.hpp @@ -24,7 +24,6 @@ #include #include -#include #include #include From 630dca2769c6b4704abd3f2162251e9a31ca32e0 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 23 Jul 2020 09:53:47 -0400 Subject: [PATCH 42/64] Update CHANGELOG message --- CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 136c3708bbd..2ffbe64885c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -48,7 +48,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 `fixed_point` Column Support (Part 1) +- PR #5704 Initial `fixed_point` Column Support - PR #5572 Add `cudf::encode` API. ## Improvements From 127524995104962afa4d181a6110eac6815ecd54 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 23 Jul 2020 10:53:19 -0400 Subject: [PATCH 43/64] Update doxygen --- cpp/include/cudf/scalar/scalar.hpp | 80 ++++++++++++++++-------------- 1 file changed, 44 insertions(+), 36 deletions(-) diff --git a/cpp/include/cudf/scalar/scalar.hpp b/cpp/include/cudf/scalar/scalar.hpp index 59b2f60c975..d647297dc02 100644 --- a/cpp/include/cudf/scalar/scalar.hpp +++ b/cpp/include/cudf/scalar/scalar.hpp @@ -93,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, @@ -163,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, @@ -181,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, @@ -216,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, @@ -234,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, @@ -266,10 +270,10 @@ class fixed_point_scalar : public detail::fixed_width_scalar { /** * @brief Construct a new fixed_point 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_point_scalar(T value, bool is_valid = true, @@ -284,6 +288,8 @@ class fixed_point_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 */ fixed_point_scalar(rmm::device_scalar&& data, bool is_valid = true, @@ -313,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, @@ -330,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, @@ -348,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, @@ -417,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, @@ -433,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, @@ -451,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, From 583b26a2d020940ed9f37d2840ba234b74a113fd Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 23 Jul 2020 13:11:23 -0400 Subject: [PATCH 44/64] Move reduction unit tests --- cpp/tests/fixed_point/fixed_point_tests.cu | 51 ------------------ cpp/tests/reductions/reduction_tests.cpp | 60 ++++++++++++++++++++++ 2 files changed, 60 insertions(+), 51 deletions(-) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 5ad6a0da076..236de535bd8 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -24,7 +24,6 @@ #include #include #include -#include #include #include #include @@ -724,54 +723,4 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointInterleave) } } -TYPED_TEST(FixedPointTestBothReps, FixedPointReductionProduct) -{ - using decimalXX = fixed_point; - - 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 = 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 decimalXX = fixed_point; - - // 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 = 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/reductions/reduction_tests.cpp b/cpp/tests/reductions/reduction_tests.cpp index 2146fec0bba..5076196db63 100644 --- a/cpp/tests/reductions/reduction_tests.cpp +++ b/cpp/tests/reductions/reduction_tests.cpp @@ -22,6 +22,7 @@ #include #include +#include #include #include @@ -926,4 +927,63 @@ TYPED_TEST(ReductionTest, UniqueCount) cudf::make_nunique_aggregation(cudf::null_policy::EXCLUDE)); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +using FixedPointTypes = ::testing::Types; +TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointReductionProduct) +{ + using namespace numeric; + using decimalXX = fixed_point; + + 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 = fixed_point; + + // 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() From 7ae74859c130000c018c712f80bde58ce2d8d372 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 23 Jul 2020 13:16:56 -0400 Subject: [PATCH 45/64] Move interleave fixed_point test --- cpp/tests/fixed_point/fixed_point_tests.cu | 22 ------------- .../reshape/interleave_columns_tests.cpp | 31 +++++++++++++++++++ 2 files changed, 31 insertions(+), 22 deletions(-) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 236de535bd8..0c4b7fbeacb 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -25,7 +25,6 @@ #include #include #include -#include #include #include #include @@ -702,25 +701,4 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointUpperBound) expect_columns_equal(*result, expect); } -TYPED_TEST(FixedPointTestBothReps, FixedPointInterleave) -{ - using decimalXX = fixed_point; - - 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 THREE = decimalXX{3, scale_type{i}}; - auto const FOUR = decimalXX{4, scale_type{i}}; - - auto const a = wrapper({ONE, THREE}); - auto const b = wrapper({TWO, FOUR}); - - auto const input = cudf::table_view{std::vector{a, b}}; - auto const expected = wrapper({ONE, TWO, THREE, FOUR}); - auto const actual = cudf::interleave_columns(input); - - expect_columns_equal(expected, actual->view()); - } -} - CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 67a683aa443..5fe53f8720e 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; @@ -343,4 +345,33 @@ TEST_F(InterleaveStringsColumnsTest, MultiColumnStringMixNullableMix) cudf::test::expect_columns_equal(*results, exp_results, true); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +using FixedPointTypes = ::testing::Types; +TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointInterleave) +{ + using namespace numeric; + using decimalXX = fixed_point; + + 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() From 8e41cd04edcc053a91356838d8061c7f2c303cd5 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 23 Jul 2020 13:20:59 -0400 Subject: [PATCH 46/64] Move search fixed_point unit tests --- cpp/tests/fixed_point/fixed_point_tests.cu | 45 ----------------- cpp/tests/search/search_test.cpp | 56 ++++++++++++++++++++++ 2 files changed, 56 insertions(+), 45 deletions(-) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 0c4b7fbeacb..aa5ff6d1749 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -28,7 +28,6 @@ #include #include #include -#include "cudf/search.hpp" #include "cudf/types.hpp" #include @@ -657,48 +656,4 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointReplace) cudf::test::expect_columns_equal(*result, expected_w); } -TYPED_TEST(FixedPointTestBothReps, FixedPointLowerBound) -{ - using decimalXX = fixed_point; - - auto vec = std::vector(1000); - std::iota(std::begin(vec), std::end(vec), decimalXX{}); - - auto const values = wrapper{decimalXX{200, scale_type{0}}, - decimalXX{400, scale_type{0}}, - decimalXX{600, scale_type{0}}, - decimalXX{800, scale_type{0}}}; - auto const expect = wrapper{200, 400, 600, 800}; - auto const 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 decimalXX = fixed_point; - - auto vec = std::vector(1000); - std::iota(std::begin(vec), std::end(vec), decimalXX{}); - - auto const values = wrapper{decimalXX{200, scale_type{0}}, - decimalXX{400, scale_type{0}}, - decimalXX{600, scale_type{0}}, - decimalXX{800, scale_type{0}}}; - auto const expect = wrapper{201, 401, 601, 801}; - auto const 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/search/search_test.cpp b/cpp/tests/search/search_test.cpp index 6e2ace07b99..3fea6ed8937 100644 --- a/cpp/tests/search/search_test.cpp +++ b/cpp/tests/search/search_test.cpp @@ -18,6 +18,7 @@ #include #include +#include #include "cudf/search.hpp" struct SearchTest : public cudf::test::BaseFixture { @@ -1815,4 +1816,59 @@ TEST_F(SearchTest, multi_contains_empty_input_set_string) expect_columns_equal(*result, expect); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +using FixedPointTypes = ::testing::Types; +TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointLowerBound) +{ + using namespace numeric; + using decimalXX = fixed_point; + + 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 = fixed_point; + + 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() From d81226d2c5e83e8c4de97ccffff8547a45d7b650 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 23 Jul 2020 13:25:11 -0400 Subject: [PATCH 47/64] Move replace fixed_point unit tests --- cpp/tests/fixed_point/fixed_point_tests.cu | 27 -------------- cpp/tests/replace/replace_tests.cpp | 42 ++++++++++++++++++++-- 2 files changed, 40 insertions(+), 29 deletions(-) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index aa5ff6d1749..365b3daa3c7 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -24,7 +24,6 @@ #include #include #include -#include #include #include #include @@ -630,30 +629,4 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointConcatentate) cudf::test::expect_columns_equal(*results, expected); } -TYPED_TEST(FixedPointTestBothReps, FixedPointReplace) -{ - using decimalXX = fixed_point; - - 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/replace/replace_tests.cpp b/cpp/tests/replace/replace_tests.cpp index ef0ac5dcebf..ad5b57b765f 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,40 @@ 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; +using FixedPointTypes = ::testing::Types; +TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointReplace) +{ + using namespace numeric; + using decimalXX = fixed_point; + + 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() From 15edb5dc98aaa72e75b7b9d1693481a8c77348f3 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 23 Jul 2020 13:29:15 -0400 Subject: [PATCH 48/64] Move concatentate fixed_point test --- cpp/tests/copying/concatenate_tests.cu | 38 +++++++++++++++++++--- cpp/tests/fixed_point/fixed_point_tests.cu | 18 ---------- 2 files changed, 34 insertions(+), 22 deletions(-) diff --git a/cpp/tests/copying/concatenate_tests.cu b/cpp/tests/copying/concatenate_tests.cu index e577857723c..d62ddcf9a69 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,31 @@ 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; +using FixedPointTypes = ::testing::Types; +TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointConcatentate) +{ + using namespace numeric; + using decimalXX = fixed_point; + + 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/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 365b3daa3c7..05d44488a69 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -611,22 +611,4 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiply) cudf::test::expect_columns_equal(expected_col, result->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointConcatentate) -{ - using decimalXX = fixed_point; - - 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); -} - CUDF_TEST_PROGRAM_MAIN() From 77d654ffa88007c323d7e77e9633f84374610d11 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 23 Jul 2020 13:35:16 -0400 Subject: [PATCH 49/64] Move binaryop fixed_point unit tests --- cpp/tests/binaryop/binop-integration-test.cpp | 71 ++++++++++++++++++- cpp/tests/fixed_point/fixed_point_tests.cu | 57 --------------- 2 files changed, 70 insertions(+), 58 deletions(-) diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 75d5f00d38d..43845f0f21e 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -18,9 +18,11 @@ */ #include -#include #include +#include +#include + namespace cudf { namespace test { namespace binop { @@ -2036,6 +2038,73 @@ 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; +using FixedPointTypes = ::testing::Types; +TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd) +{ + using namespace numeric; + using decimalXX = fixed_point; + + 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 = fixed_point; + + 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()); +} + } // namespace binop } // namespace test } // namespace cudf diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 05d44488a69..afaf22acfe5 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -20,7 +20,6 @@ #include "tests/utilities/column_utilities.hpp" #include "tests/utilities/table_utilities.hpp" -#include #include #include #include @@ -555,60 +554,4 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointSortedOrderGather) cudf::test::expect_tables_equal(sorted_table, sorted->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd) -{ - using decimalXX = fixed_point; - - 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 decimalXX = fixed_point; - - 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()); -} - CUDF_TEST_PROGRAM_MAIN() From 0c8b05bf4f070f774356e17ebaa59f91f3117686 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 23 Jul 2020 13:43:54 -0400 Subject: [PATCH 50/64] Move sort fixed_point unit test --- cpp/tests/fixed_point/fixed_point_tests.cu | 38 ---------------- cpp/tests/sort/sort_test.cpp | 51 +++++++++++++++++++--- 2 files changed, 46 insertions(+), 43 deletions(-) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index afaf22acfe5..cfd669f8246 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -15,18 +15,11 @@ */ #include -#include #include -#include "tests/utilities/column_utilities.hpp" -#include "tests/utilities/table_utilities.hpp" #include -#include #include -#include -#include #include -#include "cudf/types.hpp" #include #include @@ -523,35 +516,4 @@ TEST_F(FixedPointTest, DecimalXXThrustOnDevice) EXPECT_EQ(vec2, vec3); } -template -using wrapper = cudf::test::fixed_width_column_wrapper; - -TYPED_TEST(FixedPointTestBothReps, FixedPointSortedOrderGather) -{ - using decimalXX = fixed_point; - - 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()); -} - CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/sort/sort_test.cpp b/cpp/tests/sort/sort_test.cpp index bc4737c4cb1..96461569dc7 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,44 @@ 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; +using FixedPointTypes = ::testing::Types; +TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointSortedOrderGather) +{ + using namespace numeric; + using decimalXX = fixed_point; + + 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 From 50558f76289c8d0d37eeecb5320a9939079b613f Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 27 Jul 2020 23:42:16 -0400 Subject: [PATCH 51/64] Add simple binary op EQUAL unit test --- cpp/tests/binaryop/binop-integration-test.cpp | 31 +++++++++++++++++++ 1 file changed, 31 insertions(+) diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 43845f0f21e..6be8ddd009f 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -22,6 +22,8 @@ #include #include +#include "cudf/types.hpp" +#include "cudf/utilities/type_dispatcher.hpp" namespace cudf { namespace test { @@ -2105,6 +2107,35 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiply) cudf::test::expect_columns_equal(expected_col, result->view()); } +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimple) +{ + using namespace numeric; + using decimalXX = fixed_point; + + 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()); +} + } // namespace binop } // namespace test } // namespace cudf From 10336570cb7077459067d69878ffe3289f5481cf Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 27 Jul 2020 23:57:49 -0400 Subject: [PATCH 52/64] Add binaryop comparison fixed_point unit tests --- cpp/tests/binaryop/binop-integration-test.cpp | 41 +++++++++++++++++++ 1 file changed, 41 insertions(+) diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 6be8ddd009f..fb024f895c1 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -24,6 +24,7 @@ #include #include "cudf/types.hpp" #include "cudf/utilities/type_dispatcher.hpp" +#include "tests/utilities/column_utilities.hpp" namespace cudf { namespace test { @@ -2136,6 +2137,46 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimple) cudf::test::expect_columns_equal(expected, result->view()); } +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualLessGreater) +{ + using namespace numeric; + using decimalXX = fixed_point; + + 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 From 4f4506d236d32d42ea3193517669bea8839c4747 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 28 Jul 2020 12:35:29 -0400 Subject: [PATCH 53/64] Add scatter fixed_point unit test --- cpp/tests/copying/scatter_tests.cpp | 34 +++++++++++++++++++++++++++++ 1 file changed, 34 insertions(+) diff --git a/cpp/tests/copying/scatter_tests.cpp b/cpp/tests/copying/scatter_tests.cpp index 3990577e264..a811b8b23db 100644 --- a/cpp/tests/copying/scatter_tests.cpp +++ b/cpp/tests/copying/scatter_tests.cpp @@ -948,3 +948,37 @@ 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; +using FixedPointTypes = ::testing::Types; +TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, FixedPointScatter) +{ + using namespace numeric; + using decimalXX = fixed_point; + + 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()); +} From 4d56b6559c3fbb1a2f61743fd9502c85589687c9 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 28 Jul 2020 14:10:11 -0400 Subject: [PATCH 54/64] Add cudf::test::FixedPointTypes type list and update unit tests --- cpp/tests/binaryop/binop-integration-test.cpp | 14 +++++++------- cpp/tests/copying/concatenate_tests.cu | 7 +++---- cpp/tests/copying/scatter_tests.cpp | 7 +++---- cpp/tests/reductions/reduction_tests.cpp | 7 +++---- cpp/tests/replace/replace_tests.cpp | 7 +++---- cpp/tests/reshape/interleave_columns_tests.cpp | 5 ++--- cpp/tests/search/search_test.cpp | 8 ++++---- cpp/tests/sort/sort_test.cpp | 7 +++---- cpp/tests/utilities/type_lists.hpp | 13 +++++++++++++ 9 files changed, 41 insertions(+), 34 deletions(-) diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index fb024f895c1..e40c6b02504 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -19,6 +19,7 @@ #include #include +#include #include #include @@ -2046,14 +2047,13 @@ struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; template -using wrapper = cudf::test::fixed_width_column_wrapper; -using FixedPointTypes = ::testing::Types; -TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); +using wrapper = cudf::test::fixed_width_column_wrapper; +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; auto const sz = std::size_t{1000}; @@ -2082,7 +2082,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd) TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiply) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; auto const sz = std::size_t{1000}; @@ -2111,7 +2111,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiply) TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimple) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; auto const ONE = decimalXX{1, scale_type{0}}; auto const TWO = decimalXX{2, scale_type{0}}; @@ -2140,7 +2140,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimple) TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualLessGreater) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; auto const sz = std::size_t{1000}; diff --git a/cpp/tests/copying/concatenate_tests.cu b/cpp/tests/copying/concatenate_tests.cu index d62ddcf9a69..449c6c4558a 100644 --- a/cpp/tests/copying/concatenate_tests.cu +++ b/cpp/tests/copying/concatenate_tests.cu @@ -554,14 +554,13 @@ struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; template -using wrapper = cudf::test::fixed_width_column_wrapper; -using FixedPointTypes = ::testing::Types; -TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); +using wrapper = cudf::test::fixed_width_column_wrapper; +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, FixedPointConcatentate) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; auto vec = std::vector(1000); std::iota(std::begin(vec), std::end(vec), decimalXX{}); diff --git a/cpp/tests/copying/scatter_tests.cpp b/cpp/tests/copying/scatter_tests.cpp index a811b8b23db..a0053a7fa2c 100644 --- a/cpp/tests/copying/scatter_tests.cpp +++ b/cpp/tests/copying/scatter_tests.cpp @@ -954,14 +954,13 @@ struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; template -using wrapper = cudf::test::fixed_width_column_wrapper; -using FixedPointTypes = ::testing::Types; -TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); +using wrapper = cudf::test::fixed_width_column_wrapper; +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, FixedPointScatter) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; auto const ONE = decimalXX{1, scale_type{0}}; auto const TWO = decimalXX{2, scale_type{0}}; diff --git a/cpp/tests/reductions/reduction_tests.cpp b/cpp/tests/reductions/reduction_tests.cpp index 5383760aeac..520b1890632 100644 --- a/cpp/tests/reductions/reduction_tests.cpp +++ b/cpp/tests/reductions/reduction_tests.cpp @@ -925,13 +925,12 @@ template struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; -using FixedPointTypes = ::testing::Types; -TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, FixedPointReductionProduct) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; auto const ONE = decimalXX{1, scale_type{0}}; auto const TWO = decimalXX{2, scale_type{0}}; @@ -957,7 +956,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointReductionProduct) TYPED_TEST(FixedPointTestBothReps, FixedPointReductionSum) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; // auto const ZERO = decimalXX{0, scale_type{0}}; auto const ONE = decimalXX{1, scale_type{0}}; diff --git a/cpp/tests/replace/replace_tests.cpp b/cpp/tests/replace/replace_tests.cpp index ad5b57b765f..fd3352d6c37 100644 --- a/cpp/tests/replace/replace_tests.cpp +++ b/cpp/tests/replace/replace_tests.cpp @@ -535,14 +535,13 @@ struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; template -using wrapper = cudf::test::fixed_width_column_wrapper; -using FixedPointTypes = ::testing::Types; -TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); +using wrapper = cudf::test::fixed_width_column_wrapper; +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, FixedPointReplace) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; auto const ONE = decimalXX{1, scale_type{0}}; auto const TWO = decimalXX{2, scale_type{0}}; diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 5fe53f8720e..4497b0d3e52 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -349,13 +349,12 @@ template struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; -using FixedPointTypes = ::testing::Types; -TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, FixedPointInterleave) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; for (int i = 0; i > -4; --i) { auto const ONE = decimalXX{1, scale_type{i}}; diff --git a/cpp/tests/search/search_test.cpp b/cpp/tests/search/search_test.cpp index 3fea6ed8937..4ca8de65e81 100644 --- a/cpp/tests/search/search_test.cpp +++ b/cpp/tests/search/search_test.cpp @@ -17,6 +17,7 @@ #include #include #include +#include #include #include "cudf/search.hpp" @@ -1820,13 +1821,12 @@ template struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; -using FixedPointTypes = ::testing::Types; -TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, FixedPointLowerBound) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; auto vec = std::vector(1000); std::iota(std::begin(vec), std::end(vec), decimalXX{}); @@ -1850,7 +1850,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointLowerBound) TYPED_TEST(FixedPointTestBothReps, FixedPointUpperBound) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; auto vec = std::vector(1000); std::iota(std::begin(vec), std::end(vec), decimalXX{}); diff --git a/cpp/tests/sort/sort_test.cpp b/cpp/tests/sort/sort_test.cpp index 96461569dc7..0d68047d0d4 100644 --- a/cpp/tests/sort/sort_test.cpp +++ b/cpp/tests/sort/sort_test.cpp @@ -290,14 +290,13 @@ struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; template -using wrapper = cudf::test::fixed_width_column_wrapper; -using FixedPointTypes = ::testing::Types; -TYPED_TEST_CASE(FixedPointTestBothReps, FixedPointTypes); +using wrapper = cudf::test::fixed_width_column_wrapper; +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, FixedPointSortedOrderGather) { using namespace numeric; - using decimalXX = fixed_point; + using decimalXX = TypeParam; auto const ZERO = decimalXX{0, scale_type{0}}; auto const ONE = decimalXX{1, scale_type{0}}; diff --git a/cpp/tests/utilities/type_lists.hpp b/cpp/tests/utilities/type_lists.hpp index 60dfbe5a3c9..71356c3328e 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. From 5b8608b0cbef1863eb087d3314067822a124c4bf Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 28 Jul 2020 14:44:34 -0400 Subject: [PATCH 55/64] Add fixed_width_type_converter specialization for fixed_point --- cpp/tests/utilities/column_wrapper.hpp | 13 +++++++++++++ 1 file changed, 13 insertions(+) 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 From 30219108d7a29a37c0336aaeb2f9490d703ce4b6 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 28 Jul 2020 14:45:19 -0400 Subject: [PATCH 56/64] Add cudf::test::FixedPointTypes to scatter invalid index test --- cpp/tests/copying/scatter_tests.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/cpp/tests/copying/scatter_tests.cpp b/cpp/tests/copying/scatter_tests.cpp index a0053a7fa2c..8fd1f45663e 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 From 5ee38c00ea94c805eea0c61646056c5f8609639f Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 28 Jul 2020 23:21:03 -0400 Subject: [PATCH 57/64] Change long to uint64_t --- cpp/src/io/json/json_gpu.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/src/io/json/json_gpu.cu b/cpp/src/io/json/json_gpu.cu index 92a515c467d..1347f2ecef5 100644 --- a/cpp/src/io/json/json_gpu.cu +++ b/cpp/src/io/json/json_gpu.cu @@ -274,8 +274,8 @@ __inline__ __device__ cudf::list_view decode_value(const char *data, template <> __inline__ __device__ numeric::decimal32 decode_value(const char *data, - long start, - long end, + uint64_t start, + uint64_t end, ParseOptions const &opts) { return numeric::decimal32{}; @@ -283,8 +283,8 @@ __inline__ __device__ numeric::decimal32 decode_value(const char *data, template <> __inline__ __device__ numeric::decimal64 decode_value(const char *data, - long start, - long end, + uint64_t start, + uint64_t end, ParseOptions const &opts) { return numeric::decimal64{}; From a69dd142b791d273e1dd08382aaa2ec02ed9507c Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 29 Jul 2020 00:00:10 -0400 Subject: [PATCH 58/64] Add FixedPointTypes to FixedWidthTypes --- cpp/tests/copying/get_value_tests.cpp | 4 ++-- cpp/tests/copying/scatter_tests.cpp | 4 ++-- cpp/tests/copying/shift_tests.cpp | 4 ++-- cpp/tests/device_atomics/device_atomics_test.cu | 2 +- cpp/tests/filling/fill_tests.cpp | 2 +- cpp/tests/filling/repeat_tests.cpp | 2 +- .../grouped_rolling/grouped_rolling_test.cpp | 2 +- cpp/tests/merge/merge_string_test.cpp | 2 +- cpp/tests/merge/merge_test.cpp | 2 +- cpp/tests/partitioning/hash_partition_test.cpp | 2 +- cpp/tests/partitioning/round_robin_test.cpp | 2 +- cpp/tests/replace/clamp_test.cpp | 2 +- cpp/tests/rolling/rolling_test.cpp | 2 +- cpp/tests/scalar/factories_test.cpp | 2 +- cpp/tests/scalar/scalar_device_view_test.cu | 2 +- cpp/tests/scalar/scalar_test.cpp | 2 +- cpp/tests/transpose/transpose_test.cpp | 2 +- cpp/tests/utilities/type_lists.hpp | 16 +++++++++++++++- .../utilities_tests/column_utilities_tests.cpp | 2 +- .../utilities_tests/column_wrapper_tests.cpp | 2 +- 20 files changed, 37 insertions(+), 23 deletions(-) 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 8fd1f45663e..f6ab4817618 100644 --- a/cpp/tests/copying/scatter_tests.cpp +++ b/cpp/tests/copying/scatter_tests.cpp @@ -302,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) @@ -796,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) { 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/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/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/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/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/type_lists.hpp b/cpp/tests/utilities/type_lists.hpp index 71356c3328e..00d41e835fb 100644 --- a/cpp/tests/utilities/type_lists.hpp +++ b/cpp/tests/utilities/type_lists.hpp @@ -205,7 +205,21 @@ using FixedPointTypes = cudf::test::Types; +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) { From 5138dd30c361d45f6d35d97aef1657c7d2598d14 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 29 Jul 2020 10:39:16 -0400 Subject: [PATCH 59/64] Fix COPY_TEST failure --- cpp/tests/copying/copy_tests.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 From eafe0be99b41e48837f63c42c7fd49a14cdd67c4 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 29 Jul 2020 10:45:26 -0400 Subject: [PATCH 60/64] Fix FACTORIES_TEST failure --- cpp/tests/column/factories_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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) { From f4946465d44dd0da20059f190f20775fda2a2e03 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 29 Jul 2020 10:51:19 -0400 Subject: [PATCH 61/64] Fix GROUPBY_TEST failure --- cpp/tests/groupby/group_max_test.cpp | 2 +- cpp/tests/groupby/group_min_test.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) 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) From b8e43a3e9c7bb0ff1a92ebac31b74c3e7ee18b59 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 29 Jul 2020 19:00:54 -0400 Subject: [PATCH 62/64] Remove `is_primitive_type` trait --- .../cudf/detail/aggregation/aggregation.cuh | 35 ++++++++++--------- cpp/include/cudf/utilities/traits.hpp | 14 -------- cpp/src/copying/get_element.cu | 2 +- cpp/src/unary/cast_ops.cu | 4 ++- 4 files changed, 23 insertions(+), 32 deletions(-) diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index 40c26d2656e..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_primitive_type() 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/utilities/traits.hpp b/cpp/include/cudf/utilities/traits.hpp index 6fe8762b290..d99d5daf10a 100644 --- a/cpp/include/cudf/utilities/traits.hpp +++ b/cpp/include/cudf/utilities/traits.hpp @@ -472,20 +472,6 @@ constexpr inline bool is_fixed_width(data_type type) return cudf::type_dispatcher(type, is_fixed_width_impl{}); } -/** - * @brief Indicates whether the type `T` is a "primitive type". Currently this - * is fixed width types not including fixed point - * - * @tparam T The type to verify - * @return true `T` is a primitive type - * @return false `T` is not a primitive type - **/ -template -constexpr inline bool is_primitive_type() -{ - return cudf::is_fixed_width() && !cudf::is_fixed_point(); -} - /** * @brief Indicates whether the type `T` is a compound type. * diff --git a/cpp/src/copying/get_element.cu b/cpp/src/copying/get_element.cu index d5746188a74..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, diff --git a/cpp/src/unary/cast_ops.cu b/cpp/src/unary/cast_ops.cu index e38316d01d2..338bb481606 100644 --- a/cpp/src/unary/cast_ops.cu +++ b/cpp/src/unary/cast_ops.cu @@ -153,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) From d2eb6b6fe355412276abd30e6669e53d8eb51105 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 30 Jul 2020 00:12:06 -0400 Subject: [PATCH 63/64] Fix ROLLING_TEST failure!!! :) --- cpp/include/cudf/detail/utilities/device_operators.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cudf/detail/utilities/device_operators.cuh b/cpp/include/cudf/detail/utilities/device_operators.cuh index 299166a242b..460b0181abe 100644 --- a/cpp/include/cudf/detail/utilities/device_operators.cuh +++ b/cpp/include/cudf/detail/utilities/device_operators.cuh @@ -47,7 +47,7 @@ struct DeviceSum { template ()>* = nullptr> static constexpr T identity() { - return T{}; + return T{0}; } template ()>* = nullptr> From c2937853d973cfc74a1eb5ee6bef26c3ed193c30 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 30 Jul 2020 21:12:43 -0400 Subject: [PATCH 64/64] Add comment on why using includes --- cpp/include/cudf/fixed_point/fixed_point.hpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index f77d3fb06ae..3634d86f37a 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -18,6 +18,8 @@ #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