From dd3eb62847b763a8579c62d2e7444e4f774e75d2 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra <36027403+codereport@users.noreply.github.com> Date: Tue, 13 Oct 2020 17:45:27 -0400 Subject: [PATCH] [REVIEW] `fixed_point` Column Optimization (store `scale` in `data_type`) [skip ci] (#5861) * Add fixed_point iterator test * Add test demonstrating binary op * Replace cout with asserts * Test with fraction fixed_point values * Attempt to implement proxy reference, but thrust doesn't like it * Add hacked up transform_mutable_iterator * Cleanup dead code * Actually make mutable iterator work as input iterator * Simplify diff from transform_output_iterator * Add device implementation of fixed point iterator * Add test with iterator random access in kernel * Use thrust/transform_input_output_iterator * Replace rescale with rescaled * Add constexpr to numeric::scale_type ctor * Cleanup forward declarations * Add scale to data_type * Fixed typo in clamp file * Add make_numeric_column * factories_test.cpp remove FixedWidthTypesWithoutFixedPoint * Update CHANGELOG * Fix mockup of ColumnLike, get unit tests working, clang-format * Update CHANGELOG * Fix formatting * Add scale_type ctor to data_type * Formatting of make_fixed_width_column for readability * Fix typo * Add `cudf::` to data_type member for consistency * Update comment * Add `scale` method to data_type * Clean up column_view_printer for fixed_point * Add decimal32 specialization for .element() * Add initial `fixed_point_column_wrapper` * Add temporary change to column_view_printer * Add unit test for simple `fixed_point_column_wrapper` * Add `thrust::optional` to `data_type::_scale` * Make `fixed_point_column_wrapper` unit test for both reps * Add decimal64 specialization for .element() * Update column_view_printer for decimal64 * Small cleanup * Fix for failing unit tests * CUDF_EXPECTS data_type::_scale to be set * Add unit test for data_type::_scale not set * Remove TODO comment * data_type ctor docs and CUDF_EXPECTS * Add unit test for wrong type_id in data_type ctor * Remove `thrust::optional` from scale in data_type * Update CHANGELOG * IIFE clean up * Return fixed_point without temporary scaled_integer * Change fixed_point unit test to 0 scale * Add cudf::distance to avoid implicit casting * Temporary fix for FILLING_TEST * Add fixed_point specialization for make_elements * Fixed FIXED_POINT_TEST * Cleanup unit test * Add specialization of cudf::to_host for fixed_point * Add another specialization of make_elements * Use .empty() instead of == 0 * Fix ROLLING_TEST * CI fix * Fix CI: get_current_default_resource * Delete file * Merge mistake: host_vector should be of RepType * Use counting iterator + transform iterator * Use lambda and transform iterator * Add const * Cleanup * Fix (can't use generic lambda) * Use std::for_each & std::any_of * Fixed `RESHAPE_TEST` & `TRANSPOSE_TEST` * Fix PARTITIONING_TEST and some of COPYING_TEST * Fix ShiftTests of COPYING_TEST * Cleanup: use ternary opreator * Fix MERGE_TEST * Temporarily disable BINOP * Fix for ScatterScalar tests of COPYING_TEST * Fix DictionaryConcat tests of COPYING_TEST * Disable copy_range for fixed_point * Add specialization for copy_range fixed_point * Temporarily disable REPLACE_TEST for fixed_point * Use snake_case cleanup * Add back REPLACE_TEST for fixed_point * Use cudf::test::make_counting_transform_iterator * Fix for REPLACE_TEST * Add get_column_stored_type * Add CUDF_EXPECTS on representation type * Fix/cleanup cudf::to_host for fixed_point * Add simple fixed_point_column_wrapper test * fixed_point column_view_printer * Small cleanup * Remove [&] * Remove 10 * (from trying to break tests) * Rename _scale and adjust comment for clarity * Add column_type_id_matches_column_stored_type * Use thrust::copy_if with stencil * Add missing header * Docs and removing TODOs * Disable binop ptx with fixed_point * Add initial support for fixed_point binary operations * Enable some of the binary op tests * Refactor / cleanup * Remove rmm_log.txt * Remove unnecessary header * Use .is_empty() instead of .size() == 0 * Make changes to fixed_point_scalar * Changes to search.cu * Changes to scalar_construction_helper * Changes to null.cu * Changes to src/dictionary/search.cu * Changes to fixed_point_scalar * Remove temporary changes * Addressing PR comments * Update cpp/include/cudf/utilities/type_dispatcher.hpp Co-authored-by: Mark Harris * Update cpp/include/cudf/column/column_factories.hpp Co-authored-by: Mark Harris * Update cpp/include/cudf/column/column_factories.hpp Co-authored-by: Mark Harris * Update cpp/include/cudf/utilities/type_dispatcher.hpp Co-authored-by: Mark Harris * Fix formatting * Changes from PR 6063 and 6064 * Addressing PR comments * CI Fix * Update cpp/src/merge/merge.cu Co-authored-by: Mark Harris * Update cpp/src/merge/merge.cu Co-authored-by: Mark Harris * Update cpp/src/merge/merge.cu Co-authored-by: Mark Harris * Update cpp/src/merge/merge.cu Co-authored-by: Mark Harris * Update cpp/src/merge/merge.cu Co-authored-by: Mark Harris * Update cpp/src/merge/merge.cu Co-authored-by: Mark Harris * Auto commit fix * Replace for_each with all_of * Cleanup in types.hpp * Changes for fixed_point binary operations * fixed_point binop different scales * Enable last binop fixed_point test * Clean up tests (remove TODO) and support both decimal32/64 * Final binop cases for fixed_point * Addressing PR comments * fixed_point_scalar_device_view inheritance from scalar_device_* * Testing CI * Final TODOs (docs) * Remove redundant is_fixed_point * Fix for AST (thank you David) * Remove headers + clean up * Remove headers * fixed_point scalar factory test * scalar_test for fixed_point * Rename fn to type_id_matched_device_storage_type * Remove header * Addressing PR comments, fixed_point_column_wrapper tests * Revert type_id_matches... changes * Simplify device_storage_type_t docs * Add failing binop test * Disable `fixed_point` binaryops * Addressing PR comments * Spelling fixes Co-authored-by: Trevor Smith Co-authored-by: Mark Harris --- CHANGELOG.md | 1 + cpp/include/cudf/ast/detail/linearizer.hpp | 2 +- cpp/include/cudf/column/column.hpp | 18 +- .../cudf/column/column_device_view.cuh | 41 +- cpp/include/cudf/column/column_factories.hpp | 58 ++- cpp/include/cudf/detail/gather.cuh | 16 +- cpp/include/cudf/detail/scatter.cuh | 8 +- cpp/include/cudf/fixed_point/fixed_point.hpp | 9 +- cpp/include/cudf/scalar/scalar.hpp | 77 ++- .../cudf/scalar/scalar_device_view.cuh | 8 +- cpp/include/cudf/scalar/scalar_factories.hpp | 23 +- cpp/include/cudf/types.hpp | 72 ++- .../cudf/utilities/type_dispatcher.hpp | 36 ++ cpp/include/cudf_test/column_utilities.hpp | 33 +- cpp/include/cudf_test/column_wrapper.hpp | 137 ++++-- cpp/src/binaryop/binaryop.cpp | 454 ++++++++++++++++-- cpp/src/binaryop/compiled/binary_ops.cu | 6 +- cpp/src/binaryop/compiled/binary_ops.hpp | 2 +- cpp/src/column/column_factories.cpp | 50 +- cpp/src/copying/concatenate.cu | 6 +- cpp/src/copying/copy_range.cu | 22 + cpp/src/copying/scatter.cu | 6 +- cpp/src/copying/shift.cu | 7 +- cpp/src/copying/split.cpp | 2 +- cpp/src/dictionary/search.cu | 28 +- cpp/src/io/avro/avro.cpp | 2 +- cpp/src/io/utilities/column_buffer.hpp | 4 +- cpp/src/jit/type.cpp | 4 +- cpp/src/merge/merge.cu | 47 +- cpp/src/reductions/simple.cuh | 4 +- cpp/src/replace/clamp.cu | 2 +- cpp/src/replace/nulls.cu | 12 +- cpp/src/replace/replace.cu | 66 ++- cpp/src/reshape/interleave_columns.cu | 22 +- cpp/src/scalar/scalar_factories.cpp | 26 +- cpp/src/search/search.cu | 15 +- cpp/tests/binaryop/binop-integration-test.cpp | 259 ++++++---- cpp/tests/column/factories_test.cpp | 2 +- cpp/tests/copying/copy_range_tests.cpp | 2 +- cpp/tests/filling/repeat_tests.cpp | 14 +- cpp/tests/fixed_point/fixed_point_tests.cu | 44 ++ cpp/tests/replace/replace_tests.cpp | 11 +- .../reshape/interleave_columns_tests.cpp | 3 +- cpp/tests/scalar/factories_test.cpp | 24 +- cpp/tests/scalar/scalar_test.cpp | 15 +- cpp/tests/utilities/column_utilities.cu | 26 +- .../column_utilities_tests.cpp | 64 ++- 47 files changed, 1384 insertions(+), 406 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 0268148f63b..4611e258c80 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -52,6 +52,7 @@ - PR #6060 Add support for all types in `Series.describe` and `DataFrame.describe` - PR #6051 Add builder API for cuIO `parquet_writer_options` and `parquet_reader_options` - PR #6067 Added compute codes for aarch64 devices +- PR #5861 `fixed_point` Column Optimization (store `scale` in `data_type`) - PR #6083 Small cleanup - PR #6355 Make sure PTDS mode is compatible between libcudf and JNI - PR #6120 Consolidate functionality in NestedHostColumnVector and HostColumnVector diff --git a/cpp/include/cudf/ast/detail/linearizer.hpp b/cpp/include/cudf/ast/detail/linearizer.hpp index 25b94509831..c9f61490c5d 100644 --- a/cpp/include/cudf/ast/detail/linearizer.hpp +++ b/cpp/include/cudf/ast/detail/linearizer.hpp @@ -54,7 +54,7 @@ enum class device_data_reference_type { * by the `row_evaluator`. * */ -struct device_data_reference { +struct alignas(8) device_data_reference { device_data_reference(device_data_reference_type reference_type, cudf::data_type data_type, cudf::size_type data_index, diff --git a/cpp/include/cudf/column/column.hpp b/cpp/include/cudf/column/column.hpp index 5dac3347d63..ce0ed412b27 100644 --- a/cpp/include/cudf/column/column.hpp +++ b/cpp/include/cudf/column/column.hpp @@ -312,15 +312,15 @@ class column { operator mutable_column_view() { return this->mutable_view(); }; private: - data_type _type{type_id::EMPTY}; ///< Logical type of elements in the column - cudf::size_type _size{}; ///< The number of elements in the column - rmm::device_buffer _data{}; ///< Dense, contiguous, type erased device memory - ///< buffer containing the column elements - rmm::device_buffer _null_mask{}; ///< Bitmask used to represent null values. - ///< May be empty if `null_count() == 0` - mutable size_type _null_count{UNKNOWN_NULL_COUNT}; ///< The number of null elements - std::vector> _children{}; ///< Depending on element type, child - ///< columns may contain additional data + cudf::data_type _type{type_id::EMPTY}; ///< Logical type of elements in the column + cudf::size_type _size{}; ///< The number of elements in the column + rmm::device_buffer _data{}; ///< Dense, contiguous, type erased device memory + ///< buffer containing the column elements + rmm::device_buffer _null_mask{}; ///< Bitmask used to represent null values. + ///< May be empty if `null_count() == 0` + mutable cudf::size_type _null_count{UNKNOWN_NULL_COUNT}; ///< The number of null elements + std::vector> _children{}; ///< Depending on element type, child + ///< columns may contain additional data }; /** @} */ // end of group diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh index f3aa15842b7..5446d9b2f29 100644 --- a/cpp/include/cudf/column/column_device_view.cuh +++ b/cpp/include/cudf/column/column_device_view.cuh @@ -18,6 +18,7 @@ #include #include #include +#include #include #include #include @@ -749,6 +750,42 @@ __device__ inline dictionary32 const column_device_view::element( return dictionary32{type_dispatcher(indices.type(), index_element_fn{}, indices, index)}; } +/** + * @brief Returns a `numeric::decimal32` element at the specified index for a `fixed_point` column. + * + * If the element at the specified index is NULL, i.e., `is_null(element_index) == true`, + * then any attempt to use the result will lead to undefined behavior. + * + * @param element_index Position of the desired element + * @return numeric::decimal32 representing the element at this index + */ +template <> +__device__ inline numeric::decimal32 const column_device_view::element( + size_type element_index) const noexcept +{ + using namespace numeric; + auto const scale = scale_type{_type.scale()}; + return decimal32{scaled_integer{data()[element_index], scale}}; +} + +/** + * @brief Returns a `numeric::decimal64` element at the specified index for a `fixed_point` column. + * + * If the element at the specified index is NULL, i.e., `is_null(element_index) == true`, + * then any attempt to use the result will lead to undefined behavior. + * + * @param element_index Position of the desired element + * @return numeric::decimal64 representing the element at this index + */ +template <> +__device__ inline numeric::decimal64 const column_device_view::element( + size_type element_index) const noexcept +{ + using namespace numeric; + auto const scale = scale_type{_type.scale()}; + return decimal64{scaled_integer{data()[element_index], scale}}; +} + namespace detail { /** * @brief value accessor of column without null bitmask @@ -774,7 +811,7 @@ struct value_accessor { */ value_accessor(column_device_view const& _col) : col{_col} { - CUDF_EXPECTS(data_type(type_to_id()) == col.type(), "the data type mismatch"); + CUDF_EXPECTS(type_id_matches_device_storage_type(col.type().id()), "the data type mismatch"); } __device__ T operator()(cudf::size_type i) const { return col.element(i); } @@ -829,7 +866,7 @@ struct mutable_value_accessor { */ mutable_value_accessor(mutable_column_device_view& _col) : col{_col} { - CUDF_EXPECTS(data_type(type_to_id()) == col.type(), "the data type mismatch"); + CUDF_EXPECTS(type_id_matches_device_storage_type(col.type().id()), "the data type mismatch"); } __device__ T& operator()(cudf::size_type i) { return col.element(i); } diff --git a/cpp/include/cudf/column/column_factories.hpp b/cpp/include/cudf/column/column_factories.hpp index 88df2f97ba0..b40089f0929 100644 --- a/cpp/include/cudf/column/column_factories.hpp +++ b/cpp/include/cudf/column/column_factories.hpp @@ -88,7 +88,61 @@ std::unique_ptr make_numeric_column( cudaStream_t stream = 0, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { - CUDF_EXPECTS(is_numeric(type) || is_fixed_point(type), "Invalid, non-numeric type."); + CUDF_EXPECTS(is_numeric(type), "Invalid, non-numeric type."); + return std::make_unique(type, + size, + rmm::device_buffer{size * cudf::size_of(type), stream, mr}, + std::forward(null_mask), + null_count); +} + +/** + * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the + * specified `fixed_point` `data_type` with an optional null mask. + * + * @note The column's null count is determined by the requested null mask `state`. + * + * @throws cudf::logic_error if `type` is not a `fixed_point` type. + * + * @param[in] type The desired `fixed_point` element type. + * @param[in] size The number of elements in the column. + * @param[in] state Optional, controls allocation/initialization of the. + * column's null mask. By default, no null mask is allocated. + * @param[in] stream CUDA stream used for device memory operations and kernel launches. + * @param[in] mr Device memory resource used to allocate the returned column's device memory. + */ +std::unique_ptr make_fixed_point_column( + data_type type, + size_type size, + mask_state state = mask_state::UNALLOCATED, + cudaStream_t stream = 0, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the + * specified `fixed_point` `data_type` with a null mask. + * + * @note null_count is optional and will be computed if not provided. + * + * @throws cudf::logic_error if `type` is not a `fixed_point` type. + * + * @param[in] type The desired `fixed_point` element type. + * @param[in] size The number of elements in the column. + * @param[in] null_mask Null mask to use for this column. + * @param[in] null_count Optional number of nulls in the null_mask. + * @param[in] stream CUDA stream used for device memory operations and kernel launches. + * @param[in] mr Device memory resource used to allocate the returned column's device memory. + */ +template +std::unique_ptr make_fixed_point_column( + data_type type, + size_type size, + B&& null_mask, + size_type null_count = cudf::UNKNOWN_NULL_COUNT, + cudaStream_t stream = 0, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + CUDF_EXPECTS(is_fixed_point(type), "Invalid, non-fixed_point type."); return std::make_unique(type, size, rmm::device_buffer{size * cudf::size_of(type), stream, mr}, @@ -267,6 +321,8 @@ std::unique_ptr make_fixed_width_column( return make_timestamp_column(type, size, std::forward(null_mask), null_count, stream, mr); } else if (is_duration(type)) { return make_duration_column(type, size, std::forward(null_mask), null_count, stream, mr); + } else if (is_fixed_point(type)) { + return make_fixed_point_column(type, size, std::forward(null_mask), null_count, stream, mr); } return make_numeric_column(type, size, std::forward(null_mask), null_count, stream, mr); } diff --git a/cpp/include/cudf/detail/gather.cuh b/cpp/include/cudf/detail/gather.cuh index f0e56154771..dd6266f258b 100644 --- a/cpp/include/cudf/detail/gather.cuh +++ b/cpp/include/cudf/detail/gather.cuh @@ -44,8 +44,6 @@ #include #include -#include - namespace cudf { namespace detail { @@ -174,14 +172,16 @@ struct column_gatherer_impl { cudaStream_t stream, rmm::mr::device_memory_resource* mr) { - size_type num_destination_rows = std::distance(gather_map_begin, gather_map_end); - cudf::mask_allocation_policy policy = cudf::mask_allocation_policy::NEVER; - std::unique_ptr destination_column = - cudf::detail::allocate_like(source_column, num_destination_rows, policy, mr, stream); + auto const num_rows = cudf::distance(gather_map_begin, gather_map_end); + auto const policy = cudf::mask_allocation_policy::NEVER; + auto destination_column = + cudf::detail::allocate_like(source_column, num_rows, policy, mr, stream); + + using Type = device_storage_type_t; - gather_helper(source_column.data(), + gather_helper(source_column.data(), source_column.size(), - destination_column->mutable_view().data(), + destination_column->mutable_view().template begin(), gather_map_begin, gather_map_end, nullify_out_of_bounds, diff --git a/cpp/include/cudf/detail/scatter.cuh b/cpp/include/cudf/detail/scatter.cuh index 53de0d4b277..dd8c983d7ad 100644 --- a/cpp/include/cudf/detail/scatter.cuh +++ b/cpp/include/cudf/detail/scatter.cuh @@ -87,13 +87,15 @@ struct column_scatterer_impl { auto result = std::make_unique(target, stream, mr); auto result_view = result->mutable_view(); + using Type = device_storage_type_t; + // NOTE use source.begin + scatter rows rather than source.end in case the // scatter map is smaller than the number of source rows thrust::scatter(rmm::exec_policy(stream)->on(stream), - source.begin(), - source.begin() + std::distance(scatter_map_begin, scatter_map_end), + source.begin(), + source.begin() + cudf::distance(scatter_map_begin, scatter_map_end), scatter_map_begin, - result_view.begin()); + result_view.begin()); return result; } diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index 2f8ca6fc39c..d45ee21a250 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -39,7 +39,7 @@ namespace numeric { template struct strong_typedef { T _t; - CUDA_HOST_DEVICE_CALLABLE explicit strong_typedef(T t) : _t(t) {} + CUDA_HOST_DEVICE_CALLABLE explicit constexpr strong_typedef(T t) : _t(t) {} CUDA_HOST_DEVICE_CALLABLE operator T() const { return _t; } }; /** \endcond */ @@ -289,6 +289,8 @@ class fixed_point { scale_type _scale; public: + using rep = Rep; + /** * @brief Constructor that will perform shifting to store value appropriately * @@ -344,6 +346,11 @@ class fixed_point { return detail::shift(static_cast(_value), detail::negate(_scale)); } + CUDA_HOST_DEVICE_CALLABLE operator scaled_integer() const + { + return scaled_integer{_value, _scale}; + } + /** * @brief Explicit conversion operator to `bool` * diff --git a/cpp/include/cudf/scalar/scalar.hpp b/cpp/include/cudf/scalar/scalar.hpp index 06c5f537fba..ed3a6aebf31 100644 --- a/cpp/include/cudf/scalar/scalar.hpp +++ b/cpp/include/cudf/scalar/scalar.hpp @@ -19,6 +19,8 @@ #include #include +#include + #include #include #include @@ -259,11 +261,13 @@ class numeric_scalar : public detail::fixed_width_scalar { * @tparam T the data type of the fixed_point number */ template -class fixed_point_scalar : public detail::fixed_width_scalar { +class fixed_point_scalar : public scalar { static_assert(is_fixed_point(), "Unexpected non-fixed_point type."); public: - fixed_point_scalar() = default; + using rep_type = typename T::rep; + + fixed_point_scalar() : scalar(data_type(type_to_id())){}; ~fixed_point_scalar() = default; fixed_point_scalar(fixed_point_scalar&& other) = default; fixed_point_scalar(fixed_point_scalar const& other) = default; @@ -271,9 +275,44 @@ class fixed_point_scalar : public detail::fixed_width_scalar { fixed_point_scalar& operator=(fixed_point_scalar&& other) = delete; /** - * @brief Construct a new fixed_point scalar object + * @brief Construct a new fixed_point scalar object from already shifted value and scale * - * @param[in] value The initial value of the scalar + * @param[in] value The initial shifted value of the fixed_point scalar + * @param[in] scale The scale of the fixed_point 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(rep_type value, + numeric::scale_type scale, + bool is_valid = true, + cudaStream_t stream = 0, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) + : scalar{data_type{type_to_id(), static_cast(scale)}, is_valid, stream, mr}, + _data{value} + { + } + + /** + * @brief Construct a new fixed_point scalar object from a value and default 0-scale + * + * @param[in] value The initial value of the fixed_point 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(rep_type value, + bool is_valid = true, + cudaStream_t stream = 0, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) + : scalar{data_type{type_to_id(), 0}, is_valid, stream, mr}, _data{value} + { + } + + /** + * @brief Construct a new fixed_point scalar object from a fixed_point number + * + * @param[in] value The fixed_point number from which the fixed_point scalar will be initialized * @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 @@ -282,8 +321,11 @@ class fixed_point_scalar : public detail::fixed_width_scalar { bool is_valid = true, cudaStream_t stream = 0, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) - : detail::fixed_width_scalar(value, is_valid, stream, mr) + : scalar{data_type{type_to_id(), 0}, is_valid, stream, mr}, + _data{numeric::scaled_integer{value}.value} { + CUDF_EXPECTS(value == (T{_data.value(), numeric::scale_type{0}}), + "scale of fixed_point value should be zero"); } /** @@ -294,13 +336,34 @@ class fixed_point_scalar : public detail::fixed_width_scalar { * @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, + fixed_point_scalar(rmm::device_scalar&& data, bool is_valid = true, cudaStream_t stream = 0, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) - : detail::fixed_width_scalar(std::forward>(data), is_valid, stream, mr) + : scalar{data_type{type_to_id()}, is_valid, stream, mr}, // note that scale is ignored here + _data{std::forward>(data)} { } + + /** + * @brief Get the value of the scalar + * + * @param stream CUDA stream used for device memory operations. + */ + rep_type value(cudaStream_t stream = 0) const { return _data.value(stream); } + + /** + * @brief Returns a raw pointer to the value in device memory + */ + rep_type* data() { return _data.data(); } + + /** + * @brief Returns a const raw pointer to the value in device memory + */ + rep_type const* data() const { return _data.data(); } + + protected: + rmm::device_scalar _data{}; ///< device memory containing the value }; /** diff --git a/cpp/include/cudf/scalar/scalar_device_view.cuh b/cpp/include/cudf/scalar/scalar_device_view.cuh index 30315bb7486..a4f404b5d19 100644 --- a/cpp/include/cudf/scalar/scalar_device_view.cuh +++ b/cpp/include/cudf/scalar/scalar_device_view.cuh @@ -209,10 +209,12 @@ 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 { +class fixed_point_scalar_device_view : public detail::scalar_device_view_base { public: - fixed_point_scalar_device_view(data_type type, T* data, bool* is_valid) - : detail::fixed_width_scalar_device_view(type, data, is_valid) + using rep_type = typename T::rep; + + fixed_point_scalar_device_view(data_type type, rep_type* data, bool* is_valid) + : detail::scalar_device_view_base(type, is_valid) { } }; diff --git a/cpp/include/cudf/scalar/scalar_factories.hpp b/cpp/include/cudf/scalar/scalar_factories.hpp index d5372df1a36..feade65f31a 100644 --- a/cpp/include/cudf/scalar/scalar_factories.hpp +++ b/cpp/include/cudf/scalar/scalar_factories.hpp @@ -15,10 +15,7 @@ */ #pragma once -#include -#include -#include -#include "scalar.hpp" +#include namespace cudf { /** @@ -134,5 +131,23 @@ std::unique_ptr make_fixed_width_scalar( return std::make_unique>(value, true, stream, mr); } +/** + * @brief Construct scalar using the given value of fixed_point type + * + * @tparam T Datatype of the value to be represented by the scalar + * @param value The value to store in the scalar object + * @param stream CUDA stream used for device memory operations. + * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool. + */ +template +std::unique_ptr make_fixed_point_scalar( + typename T::rep value, + numeric::scale_type scale, + cudaStream_t stream = 0, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + return std::make_unique>(value, scale, true, stream, mr); +} + /** @} */ // end of group } // namespace cudf diff --git a/cpp/include/cudf/types.hpp b/cpp/include/cudf/types.hpp index 8b7b885d119..4107dc84909 100644 --- a/cpp/include/cudf/types.hpp +++ b/cpp/include/cudf/types.hpp @@ -24,8 +24,12 @@ #define CUDA_DEVICE_CALLABLE inline #endif +#include // TODO no idea why this is needed ¯\_(ツ)_/¯ + +#include #include #include +#include /** * @file @@ -62,31 +66,21 @@ class list_view; class struct_view; class scalar; -template -class numeric_scalar; - -template -class fixed_point_scalar; +// clang-format off +class list_scalar; class string_scalar; -template -class timestamp_scalar; -template -class duration_scalar; - -template -class numeric_scalar_device_view; - -template -class fixed_point_scalar_device_view; +template class numeric_scalar; +template class fixed_point_scalar; +template class timestamp_scalar; +template class duration_scalar; class string_scalar_device_view; -template -class timestamp_scalar_device_view; -template -class duration_scalar_device_view; - -class list_scalar; +template class numeric_scalar_device_view; +template class fixed_point_scalar_device_view; +template class timestamp_scalar_device_view; +template class duration_scalar_device_view; +// clang-format on class struct_scalar; @@ -104,6 +98,20 @@ using size_type = int32_t; using bitmask_type = uint32_t; using valid_type = uint8_t; +/** + * @brief Similar to `std::distance` but returns `cudf::size_type` and performs `static_cast` + * + * @tparam T Iterator type + * @param f "first" iterator + * @param l "last" iterator + * @return size_type The distance between first and last + */ +template +size_type distance(T f, T l) +{ + return static_cast(std::distance(f, l)); +} + /** * @brief Indicates an unknown null count. * @@ -246,15 +254,33 @@ class data_type { **/ explicit constexpr data_type(type_id id) : _id{id} {} + /** + * @brief Construct a new `data_type` object for `numeric::fixed_point` + * + * @param id The `fixed_point`'s identifier + * @param scale The `fixed_point`'s scale (see `fixed_point::_scale`) + **/ + explicit data_type(type_id id, int32_t scale) : _id{id}, _fixed_point_scale{scale} + { + assert(id == type_id::DECIMAL32 || id == type_id::DECIMAL64); + } + /** * @brief Returns the type identifier **/ CUDA_HOST_DEVICE_CALLABLE type_id id() const noexcept { return _id; } + /** + * @brief Returns the scale (for fixed_point types) + **/ + CUDA_HOST_DEVICE_CALLABLE int32_t scale() const noexcept { return _fixed_point_scale; } + private: type_id _id{type_id::EMPTY}; - // Store additional type specific metadata, timezone, decimal precision and - // scale, etc. + + // Below is additional type specific metadata. Currently, only _fixed_point_scale is stored. + + int32_t _fixed_point_scale{}; // numeric::scale_type not available here, use int32_t }; /** diff --git a/cpp/include/cudf/utilities/type_dispatcher.hpp b/cpp/include/cudf/utilities/type_dispatcher.hpp index c3910dbce24..6a8e92c4052 100644 --- a/cpp/include/cudf/utilities/type_dispatcher.hpp +++ b/cpp/include/cudf/utilities/type_dispatcher.hpp @@ -82,6 +82,42 @@ struct id_to_type_impl { template using id_to_type = typename id_to_type_impl::type; +/** + * @brief "Returns" the corresponding type that is stored on the device when using `cudf::column` + * + * For `decimal32`, the storage type is an `int32_t`. + * For `decimal64`, the storage type is an `int64_t`. + * + * Use this "type function" with the `using` type alias: + * @code + * using Type = device_storage_type_t; + * @endcode + * + * @tparam T The literal type that is stored on the host + */ +// clang-format off +template +using device_storage_type_t = + std::conditional_t::value, int32_t, + std::conditional_t::value, int64_t, T>>; +// clang-format on + +/** + * @brief Checks if `fixed_point`-like types have template type `T` matching the column's + * stored type id + * + * @tparam T The type that is stored on the device + * @param id The `data_type::id` of the column + * @return true If T matches the stored column type id + * @return false If T does not match the stored column type id + */ +template +bool type_id_matches_device_storage_type(type_id const& id) +{ + return (id == type_id::DECIMAL32 && std::is_same::value) || + (id == type_id::DECIMAL64 && std::is_same::value) || id == type_to_id(); +} + /** * @brief Macro used to define a mapping between a concrete C++ type and a *`cudf::type_id` enum. diff --git a/cpp/include/cudf_test/column_utilities.hpp b/cpp/include/cudf_test/column_utilities.hpp index b87c872907c..62eb1c084b8 100644 --- a/cpp/include/cudf_test/column_utilities.hpp +++ b/cpp/include/cudf_test/column_utilities.hpp @@ -22,6 +22,8 @@ #include #include #include +#include "cudf/utilities/traits.hpp" +#include "thrust/iterator/transform_iterator.h" namespace cudf { namespace test { @@ -146,7 +148,7 @@ bool validate_host_masks(std::vector const& expected_mask, * @return std::pair, std::vector> first is the * `column_view`'s data, and second is the column's bitmask. */ -template +template ()>* = nullptr> std::pair, std::vector> to_host(column_view c) { thrust::host_vector host_data(c.size()); @@ -154,6 +156,35 @@ std::pair, std::vector> to_host(column_view return {host_data, bitmask_to_host(c)}; } +/** + * @brief Copies the data and bitmask of a `column_view` to the host. + * + * This is the specialization for `fixed_point` that performs construction of a `fixed_point` from + * the underlying `rep` type that is stored on the device. + * + * @tparam T The data type of the elements of the `column_view` + * @param c the `column_view` to copy from + * @return std::pair, std::vector> first is the + * `column_view`'s data, and second is the column's bitmask. + */ +template ()>* = nullptr> +std::pair, std::vector> to_host(column_view c) +{ + using namespace numeric; + using Rep = typename T::rep; + + auto host_rep_types = thrust::host_vector(c.size()); + + CUDA_TRY(cudaMemcpy( + host_rep_types.data(), c.begin(), c.size() * sizeof(Rep), cudaMemcpyDeviceToHost)); + + auto to_fp = [&](Rep val) { return T{scaled_integer{val, scale_type{c.type().scale()}}}; }; + auto begin = thrust::make_transform_iterator(std::cbegin(host_rep_types), to_fp); + auto const host_fixed_points = thrust::host_vector(begin, begin + c.size()); + + return {host_fixed_points, bitmask_to_host(c)}; +} + /** * @brief Copies the data and bitmask of a `column_view` of strings * column to the host. diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp index 33f6342921b..43b5066ffbe 100644 --- a/cpp/include/cudf_test/column_wrapper.hpp +++ b/cpp/include/cudf_test/column_wrapper.hpp @@ -27,6 +27,7 @@ #include #include #include + #include #include @@ -145,28 +146,81 @@ struct fixed_width_type_converter { }; /** - * @brief Creates a `device_buffer` containing the elements in the range - * `[begin,end)`. + * @brief Creates a `device_buffer` containing the elements in the range `[begin,end)`. * - * @tparam ElementTo The type of element that is being created - * @tparam ElementFrom The type of element that is used to create elements of type `ElementTo` + * @tparam ElementTo The element type that is being created (non-`fixed_point`) + * @tparam ElementFrom The element type used to create elements of type `ElementTo` * @tparam InputIterator Iterator type for `begin` and `end` * @param begin Beginning of the sequence of elements * @param end End of the sequence of elements - * @return rmm::device_buffer Buffer containing all elements in the range - *`[begin,end)` + * @return rmm::device_buffer Buffer containing all elements in the range `[begin,end)` **/ -template +template ()>* = nullptr> rmm::device_buffer make_elements(InputIterator begin, InputIterator end) { static_assert(cudf::is_fixed_width(), "Unexpected non-fixed width type."); auto transformer = fixed_width_type_converter{}; auto transform_begin = thrust::make_transform_iterator(begin, transformer); - auto const size = std::distance(begin, end); + auto const size = cudf::distance(begin, end); auto const elements = thrust::host_vector(transform_begin, transform_begin + size); return rmm::device_buffer{elements.data(), size * sizeof(ElementTo)}; } +/** + * @brief Creates a `device_buffer` containing the elements in the range `[begin,end)`. + * + * @tparam ElementTo The element type that is being created (`fixed_point` specialization) + * @tparam ElementFrom The element type used to create elements of type `ElementTo` + * (non-`fixed-point`) + * @tparam InputIterator Iterator type for `begin` and `end` + * @param begin Beginning of the sequence of elements + * @param end End of the sequence of elements + * @return rmm::device_buffer Buffer containing all elements in the range `[begin,end)` + **/ +template () and + cudf::is_fixed_point()>* = nullptr> +rmm::device_buffer make_elements(InputIterator begin, InputIterator end) +{ + using RepType = typename ElementTo::rep; + auto transformer = fixed_width_type_converter{}; + auto transform_begin = thrust::make_transform_iterator(begin, transformer); + auto const size = cudf::distance(begin, end); + auto const elements = thrust::host_vector(transform_begin, transform_begin + size); + return rmm::device_buffer{elements.data(), size * sizeof(RepType)}; +} + +/** + * @brief Creates a `device_buffer` containing the elements in the range `[begin,end)`. + * + * @tparam ElementTo The element type that is being created (`fixed_point` specialization) + * @tparam ElementFrom The element type used to create elements of type `ElementTo` (`fixed_point`) + * @tparam InputIterator Iterator type for `begin` and `end` + * @param begin Beginning of the sequence of elements + * @param end End of the sequence of elements + * @return rmm::device_buffer Buffer containing all elements in the range `[begin,end)` + **/ +template () and + cudf::is_fixed_point()>* = nullptr> +rmm::device_buffer make_elements(InputIterator begin, InputIterator end) +{ + using namespace numeric; + using RepType = typename ElementTo::rep; + auto to_rep = [](ElementTo fp) { return static_cast>(fp).value; }; + auto transformer_begin = thrust::make_transform_iterator(begin, to_rep); + auto const size = cudf::distance(begin, end); + auto const elements = thrust::host_vector(transformer_begin, transformer_begin + size); + return rmm::device_buffer{elements.data(), size * sizeof(RepType)}; +} + /** * @brief Create a `std::vector` containing a validity indicator bitmask using * the range `[begin,end)` interpreted as booleans to indicate the state of @@ -183,12 +237,13 @@ rmm::device_buffer make_elements(InputIterator begin, InputIterator end) template std::vector make_null_mask_vector(ValidityIterator begin, ValidityIterator end) { - cudf::size_type size = std::distance(begin, end); - auto num_words = cudf::bitmask_allocation_size_bytes(size) / sizeof(bitmask_type); - std::vector null_mask(num_words, 0); - for (auto i = 0; i < size; ++i) { - if (*(begin + i)) { set_bit_unsafe(null_mask.data(), i); } - } + auto const size = cudf::distance(begin, end); + auto const num_words = cudf::bitmask_allocation_size_bytes(size) / sizeof(bitmask_type); + + auto null_mask = std::vector(num_words, 0); + for (auto i = 0; i < size; ++i) + if (*(begin + i)) set_bit_unsafe(null_mask.data(), i); + return null_mask; } @@ -285,7 +340,7 @@ class fixed_width_column_wrapper : public detail::column_wrapper { template fixed_width_column_wrapper(InputIterator begin, InputIterator end) : column_wrapper{} { - cudf::size_type size = std::distance(begin, end); + auto const size = cudf::distance(begin, end); wrapped.reset(new cudf::column{cudf::data_type{cudf::type_to_id()}, size, detail::make_elements(begin, end)}); @@ -318,8 +373,7 @@ class fixed_width_column_wrapper : public detail::column_wrapper { fixed_width_column_wrapper(InputIterator begin, InputIterator end, ValidityIterator v) : column_wrapper{} { - cudf::size_type size = std::distance(begin, end); - + auto const size = cudf::distance(begin, end); wrapped.reset(new cudf::column{cudf::data_type{cudf::type_to_id()}, size, detail::make_elements(begin, end), @@ -447,6 +501,33 @@ class fixed_width_column_wrapper : public detail::column_wrapper { } }; +template +class fixed_point_column_wrapper : public detail::column_wrapper { + public: + template + fixed_point_column_wrapper(FixedPointRepIterator begin, + FixedPointRepIterator end, + numeric::scale_type scale) + : column_wrapper{} + { + CUDF_EXPECTS(numeric::is_supported_representation_type(), "not valid representation type"); + + auto const size = cudf::distance(begin, end); + auto const elements = thrust::host_vector(begin, end); + auto const is_decimal32 = std::is_same::value; + auto const id = is_decimal32 ? type_id::DECIMAL32 : type_id::DECIMAL64; + auto const data_type = cudf::data_type{id, static_cast(scale)}; + + wrapped.reset( + new cudf::column{data_type, size, rmm::device_buffer{elements.data(), size * sizeof(Rep)}}); + } + + fixed_point_column_wrapper(std::initializer_list values, numeric::scale_type scale) + : fixed_point_column_wrapper(std::cbegin(values), std::cend(values), scale) + { + } +}; + /** * @brief `column_wrapper` derived class for wrapping columns of strings. **/ @@ -926,8 +1007,8 @@ class lists_column_wrapper : public detail::column_wrapper { // generate offsets size_type count = 0; std::vector offsetv; - std::transform(cols.begin(), - cols.end(), + std::transform(cols.cbegin(), + cols.cend(), valids, std::back_inserter(offsetv), [&](cudf::column_view const& col, bool valid) { @@ -943,10 +1024,12 @@ class lists_column_wrapper : public detail::column_wrapper { // concatenate them together, skipping children that are null. std::vector children; - for (size_t idx = 0; idx < cols.size(); idx++) { - if (!valids[idx]) { continue; } - children.push_back(cols[idx]); - } + thrust::copy_if(std::cbegin(cols), + std::cend(cols), + valids, // stencil + std::back_inserter(children), + thrust::identity{}); + auto data = children.empty() ? cudf::empty_like(expected_hierarchy) : concatenate(children); // increment depth @@ -1074,9 +1157,9 @@ class lists_column_wrapper : public detail::column_wrapper { // "a List>> that's empty at the top level" // // { {{{1, 2, 3}}}, {4, 5, 6} } - // In this case, row 1 is a a concrete List with actual values. There - // is no way to rectify the differences so we will treat it as a true - // column mismatch. + // In this case, row 1 is a a concrete List with actual values. + // There is no way to rectify the differences so we will treat it as a + // true column mismatch. CUDF_EXPECTS(l.wrapped->size() == 0, "Mismatch in column types!"); stubs.push_back(empty_like(expected_hierarchy)); } else { @@ -1227,7 +1310,7 @@ class structs_column_wrapper : public detail::column_wrapper { template void init(std::vector>&& child_columns, V validity_iterator) { - size_type num_rows = child_columns.empty() ? 0 : child_columns[0]->size(); + size_type const num_rows = child_columns.empty() ? 0 : child_columns[0]->size(); CUDF_EXPECTS(std::all_of(child_columns.begin(), child_columns.end(), diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 0cdd2db67cb..cc49a75481b 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -22,6 +22,7 @@ #include #include #include +#include #include #include #include @@ -34,6 +35,9 @@ #include // replace eventually #include "compiled/binary_ops.hpp" +#include "cudf/binaryop.hpp" +#include "cudf/fixed_point/fixed_point.hpp" +#include "cudf/types.hpp" #include #include @@ -55,11 +59,11 @@ rmm::device_buffer scalar_col_valid_mask_and(column_view const& col, cudaStream_t stream, rmm::mr::device_memory_resource* mr) { - if (col.size() == 0) { return rmm::device_buffer{0, stream, mr}; } + if (col.is_empty()) return rmm::device_buffer{0, stream, mr}; if (not s.is_valid()) { return create_null_mask(col.size(), mask_state::ALL_NULL, stream, mr); - } else if (s.is_valid() && col.nullable()) { + } else if (s.is_valid() and col.nullable()) { return copy_bitmask(col, stream, mr); } else { return rmm::device_buffer{0, stream, mr}; @@ -102,7 +106,7 @@ void binary_operation(mutable_column_view& out, binary_operator op, cudaStream_t stream) { - if (null_using_binop(op)) { + if (is_null_dependent(op)) { cudf::jit::launcher( hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream) .set_kernel_inst("kernel_v_s_with_validity", // name of the kernel we are @@ -141,7 +145,7 @@ void binary_operation(mutable_column_view& out, binary_operator op, cudaStream_t stream) { - if (null_using_binop(op)) { + if (is_null_dependent(op)) { cudf::jit::launcher( hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream) .set_kernel_inst("kernel_v_s_with_validity", // name of the kernel we are @@ -180,7 +184,7 @@ void binary_operation(mutable_column_view& out, binary_operator op, cudaStream_t stream) { - if (null_using_binop(op)) { + if (is_null_dependent(op)) { cudf::jit::launcher( hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream) .set_kernel_inst("kernel_v_v_with_validity", // name of the kernel we are @@ -247,6 +251,383 @@ void binary_operation(mutable_column_view& out, namespace detail { +// There are 3 overloads of each of the following functions: +// - `make_fixed_width_column_for_output` +// - `fixed_point_binary_operation` +// - `binary_operation` + +// The overloads are overloaded on the first two parameters of each function: +// - scalar const& lhs, column_view const& rhs, +// - column_view const& lhs, scalar const& rhs +// - column_view const& lhs, column_view const& rhs, + +/** + * @brief Helper function for making output column for binary operation + * + * @param lhs Left-hand side `scalar` used in the binary operation + * @param rhs Right-hand side `column_view` used in the binary operation + * @param op `binary_operator` to be used to combine `lhs` and `rhs` + * @param output_type `data_type` of the output column + * @param mr Device memory resource to use for device memory allocation + * @param stream CUDA stream used for device memory operations + * @return std::unique_ptr Output column used for binary operation + */ +std::unique_ptr make_fixed_width_column_for_output(scalar const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) +{ + if (binops::is_null_dependent(op)) { + return make_fixed_width_column(output_type, rhs.size(), mask_state::ALL_VALID, stream, mr); + } else { + auto new_mask = binops::detail::scalar_col_valid_mask_and(rhs, lhs, stream, mr); + return make_fixed_width_column( + output_type, rhs.size(), std::move(new_mask), cudf::UNKNOWN_NULL_COUNT, stream, mr); + } +}; + +/** + * @brief Helper function for making output column for binary operation + * + * @param lhs Left-hand side `column_view` used in the binary operation + * @param rhs Right-hand side `scalar` used in the binary operation + * @param op `binary_operator` to be used to combine `lhs` and `rhs` + * @param output_type `data_type` of the output column + * @param mr Device memory resource to use for device memory allocation + * @param stream CUDA stream used for device memory operations + * @return std::unique_ptr Output column used for binary operation + */ +std::unique_ptr make_fixed_width_column_for_output(column_view const& lhs, + scalar const& rhs, + binary_operator op, + data_type output_type, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) +{ + if (binops::is_null_dependent(op)) { + return make_fixed_width_column(output_type, lhs.size(), mask_state::ALL_VALID, stream, mr); + } else { + auto new_mask = binops::detail::scalar_col_valid_mask_and(lhs, rhs, stream, mr); + return make_fixed_width_column( + output_type, lhs.size(), std::move(new_mask), cudf::UNKNOWN_NULL_COUNT, stream, mr); + } +}; + +/** + * @brief Helper function for making output column for binary operation + * + * @param lhs Left-hand side `column_view` used in the binary operation + * @param rhs Right-hand side `column_view` used in the binary operation + * @param op `binary_operator` to be used to combine `lhs` and `rhs` + * @param output_type `data_type` of the output column + * @param mr Device memory resource to use for device memory allocation + * @param stream CUDA stream used for device memory operations + * @return std::unique_ptr Output column used for binary operation + */ +std::unique_ptr make_fixed_width_column_for_output(column_view const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) +{ + if (binops::is_null_dependent(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); + } +}; + +/** + * @brief Returns `true` if `binary_operator` `op` is a basic arithmetic binary operation + */ +bool is_basic_arithmetic_binop(binary_operator op) +{ + return op == binary_operator::ADD or ///< operator + + op == binary_operator::SUB or ///< operator - + op == binary_operator::MUL or ///< operator * + op == binary_operator::DIV; ///< operator / using common type of lhs and rhs +} + +/** + * @brief Returns `true` if `binary_operator` `op` is a comparison binary operation + */ +bool is_comparison_binop(binary_operator op) +{ + return op == binary_operator::EQUAL or ///< operator == + op == binary_operator::NOT_EQUAL or ///< operator != + op == binary_operator::LESS or ///< operator < + op == binary_operator::GREATER or ///< operator > + op == binary_operator::LESS_EQUAL or ///< operator <= + op == binary_operator::GREATER_EQUAL; ///< operator >= +} + +/** + * @brief Returns `true` if `binary_operator` `op` is supported by `fixed_point` + */ +bool is_supported_fixed_point_binop(binary_operator op) +{ + // TODO in a follow up PR + return false; // is_basic_arithmetic_binop(op) or is_comparison_binop(op); +} + +/** + * @brief Computes the scale for a `fixed_point` number based on given binary operator `op` + * + * @param op The binary_operator used for two `fixed_point` numbers + * @param left_scale Scale of left `fixed_point` number + * @param right_scale Scale of right `fixed_point` number + * @return int32_t The resulting `scale` of the computed `fixed_point` number + */ +int32_t compute_scale_for_binop(binary_operator op, int32_t left_scale, int32_t right_scale) +{ + if (op == binary_operator::MUL) return left_scale + right_scale; + if (op == binary_operator::DIV) return left_scale - right_scale; + return std::min(left_scale, right_scale); +} + +/** + * @brief Helper predicate function that identifies if `op` requires scales to be the same + * + * @param op `binary_operator` + * @return true `op` requires scales of lhs and rhs to be the same + * @return false `op` does not require scales of lhs and rhs to be the same + */ +bool is_same_scale_necessary(binary_operator op) +{ + return op != binary_operator::MUL && op != binary_operator::DIV; +} + +/** + * @brief Function to compute binary operation of one `column_view` and one `scalar` + * + * @param lhs Left-hand side `scalar` used in the binary operation + * @param rhs Right-hand side `column_view` used in the binary operation + * @param op `binary_operator` to be used to combine `lhs` and `rhs` + * @param mr Device memory resource to use for device memory allocation + * @param stream CUDA stream used for device memory operations + * @return std::unique_ptr Resulting output column from the binary operation + */ +std::unique_ptr fixed_point_binary_operation(scalar const& lhs, + column_view const& rhs, + binary_operator op, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) +{ + using namespace numeric; + + CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), + "Both columns must be of the same fixed_point type"); + + auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); + auto const output_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} // + : data_type{lhs.type().id(), scale}; + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, mr, stream); + + if (rhs.is_empty()) return out; + + auto out_view = out->mutable_view(); + + if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { + // Adjust scalar/column so they have they same scale + if (rhs.type().scale() < lhs.type().scale()) { + auto const diff = lhs.type().scale() - rhs.type().scale(); + if (lhs.type().id() == type_id::DECIMAL32) { + auto const factor = numeric::detail::ipow(diff); + auto const val = static_cast const&>(lhs).value(); + auto const scale = scale_type{rhs.type().scale()}; + auto const scalar = make_fixed_point_scalar(val * factor, scale); + binops::jit::binary_operation(out_view, *scalar, rhs, op, stream); + return out; + } else { + CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); + auto const factor = numeric::detail::ipow(diff); + auto const val = static_cast const&>(lhs).value(); + auto const scale = scale_type{rhs.type().scale()}; + auto const scalar = make_fixed_point_scalar(val * factor, scale); + binops::jit::binary_operation(out_view, *scalar, rhs, op, stream); + return out; + } + } else { + auto const diff = rhs.type().scale() - lhs.type().scale(); + auto const result = [&] { + if (lhs.type().id() == type_id::DECIMAL32) { + auto const factor = numeric::detail::ipow(diff); + auto const scalar = + make_fixed_point_scalar(factor, scale_type{rhs.type().scale()}); + return cudf::binary_operation(*scalar, rhs, binary_operator::MUL, lhs.type()); + } else { + CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); + auto const factor = numeric::detail::ipow(diff); + auto const scalar = + make_fixed_point_scalar(factor, scale_type{rhs.type().scale()}); + return cudf::binary_operation(*scalar, rhs, binary_operator::MUL, lhs.type()); + } + }(); + binops::jit::binary_operation(out_view, lhs, result->view(), op, stream); + return out; + } + } else { + binops::jit::binary_operation(out_view, lhs, rhs, op, stream); + return out; + } +} + +/** + * @brief Function to compute binary operation of one `column_view` and one `scalar` + * + * @param lhs Left-hand side `column_view` used in the binary operation + * @param rhs Right-hand side `scalar` used in the binary operation + * @param op `binary_operator` to be used to combine `lhs` and `rhs` + * @param mr Device memory resource to use for device memory allocation + * @param stream CUDA stream used for device memory operations + * @return std::unique_ptr Resulting output column from the binary operation + */ +std::unique_ptr fixed_point_binary_operation(column_view const& lhs, + scalar const& rhs, + binary_operator op, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) +{ + using namespace numeric; + + CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), + "Both columns must be of the same fixed_point type"); + + auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); + auto const output_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} // + : data_type{lhs.type().id(), scale}; + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, mr, stream); + + if (lhs.is_empty()) return out; + + auto out_view = out->mutable_view(); + + if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { + // Adjust scalar/column so they have they same scale + if (rhs.type().scale() > lhs.type().scale()) { + auto const diff = rhs.type().scale() - lhs.type().scale(); + if (rhs.type().id() == type_id::DECIMAL32) { + auto const factor = numeric::detail::ipow(diff); + auto const val = static_cast const&>(rhs).value(); + auto const scale = scale_type{lhs.type().scale()}; + auto const scalar = make_fixed_point_scalar(val * factor, scale); + binops::jit::binary_operation(out_view, lhs, *scalar, op, stream); + return out; + } else { + CUDF_EXPECTS(rhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); + auto const factor = numeric::detail::ipow(diff); + auto const val = static_cast const&>(rhs).value(); + auto const scale = scale_type{rhs.type().scale()}; + auto const scalar = make_fixed_point_scalar(val * factor, scale); + binops::jit::binary_operation(out_view, lhs, *scalar, op, stream); + return out; + } + } else { + auto const diff = lhs.type().scale() - rhs.type().scale(); + auto const result = [&] { + if (rhs.type().id() == type_id::DECIMAL32) { + auto const factor = numeric::detail::ipow(diff); + auto const scalar = + make_fixed_point_scalar(factor, scale_type{lhs.type().scale()}); + return cudf::binary_operation(*scalar, lhs, binary_operator::MUL, rhs.type()); + } else { + CUDF_EXPECTS(rhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); + auto const factor = numeric::detail::ipow(diff); + auto const scalar = + make_fixed_point_scalar(factor, scale_type{lhs.type().scale()}); + return cudf::binary_operation(*scalar, lhs, binary_operator::MUL, rhs.type()); + } + }(); + binops::jit::binary_operation(out_view, result->view(), rhs, op, stream); + return out; + } + } else { + binops::jit::binary_operation(out_view, lhs, rhs, op, stream); + return out; + } +} + +/** + * @brief Function to compute binary operation of two `column_view`s + * + * @param lhs Left-hand side `column_view` used in the binary operation + * @param rhs Right-hand side `column_view` used in the binary operation + * @param op `binary_operator` to be used to combine `lhs` and `rhs` + * @param mr Device memory resource to use for device memory allocation + * @param stream CUDA stream used for device memory operations + * @return std::unique_ptr Resulting output column from the binary operation + */ +std::unique_ptr fixed_point_binary_operation(column_view const& lhs, + column_view const& rhs, + binary_operator op, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) +{ + using namespace numeric; + + CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation"); + CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), + "Both columns must be of the same fixed_point type"); + + auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); + auto const output_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} // + : data_type{lhs.type().id(), scale}; + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, mr, stream); + + if (lhs.is_empty() or rhs.is_empty()) return out; + + auto out_view = out->mutable_view(); + + if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { + // Adjust columns so they have they same scale + if (rhs.type().scale() < lhs.type().scale()) { + auto const diff = lhs.type().scale() - rhs.type().scale(); + auto const result = [&] { + if (lhs.type().id() == type_id::DECIMAL32) { + auto const factor = numeric::detail::ipow(diff); + auto const scalar = + make_fixed_point_scalar(factor, scale_type{lhs.type().scale()}); + return cudf::binary_operation(*scalar, lhs, binary_operator::MUL, rhs.type()); + } else { + CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); + auto const factor = numeric::detail::ipow(diff); + auto const scalar = + make_fixed_point_scalar(factor, scale_type{lhs.type().scale()}); + return cudf::binary_operation(*scalar, lhs, binary_operator::MUL, rhs.type()); + } + }(); + binops::jit::binary_operation(out_view, result->view(), rhs, op, stream); + return out; + } else { + auto const diff = rhs.type().scale() - lhs.type().scale(); + auto const result = [&] { + if (lhs.type().id() == type_id::DECIMAL32) { + auto const factor = numeric::detail::ipow(diff); + auto const scalar = + make_fixed_point_scalar(factor, scale_type{rhs.type().scale()}); + return cudf::binary_operation(*scalar, rhs, binary_operator::MUL, lhs.type()); + } else { + CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); + auto const factor = numeric::detail::ipow(diff); + auto const scalar = + make_fixed_point_scalar(factor, scale_type{rhs.type().scale()}); + return cudf::binary_operation(*scalar, rhs, binary_operator::MUL, lhs.type()); + } + }(); + binops::jit::binary_operation(out_view, lhs, result->view(), op, stream); + return out; + } + } else { + binops::jit::binary_operation(out_view, lhs, rhs, op, stream); + return out; + } +} + std::unique_ptr binary_operation(scalar const& lhs, column_view const& rhs, binary_operator op, @@ -254,26 +635,20 @@ std::unique_ptr binary_operation(scalar const& lhs, rmm::mr::device_memory_resource* mr, cudaStream_t stream) { - if ((lhs.type().id() == type_id::STRING) && (rhs.type().id() == type_id::STRING)) { + if (lhs.type().id() == type_id::STRING and rhs.type().id() == type_id::STRING) return binops::compiled::binary_operation(lhs, rhs, op, output_type, mr, stream); - } + + if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) + return fixed_point_binary_operation(lhs, rhs, op, mr, stream); // Check for datatype CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype"); - 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 = binops::detail::scalar_col_valid_mask_and(rhs, lhs, stream, mr); - out = make_fixed_width_column( - output_type, rhs.size(), std::move(new_mask), cudf::UNKNOWN_NULL_COUNT, stream, mr); - } + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, mr, stream); - if (rhs.size() == 0) { return out; } + if (rhs.is_empty()) return out; auto out_view = out->mutable_view(); binops::jit::binary_operation(out_view, lhs, rhs, op, stream); @@ -287,26 +662,20 @@ std::unique_ptr binary_operation(column_view const& lhs, rmm::mr::device_memory_resource* mr, cudaStream_t stream) { - if ((lhs.type().id() == type_id::STRING) && (rhs.type().id() == type_id::STRING)) { + if (lhs.type().id() == type_id::STRING and rhs.type().id() == type_id::STRING) return binops::compiled::binary_operation(lhs, rhs, op, output_type, mr, stream); - } + + if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) + return fixed_point_binary_operation(lhs, rhs, op, mr, stream); // Check for datatype CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype"); - 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, lhs.size(), mask_state::ALL_VALID, stream, mr); - } else { - auto new_mask = binops::detail::scalar_col_valid_mask_and(lhs, rhs, stream, mr); - out = make_fixed_width_column( - output_type, lhs.size(), std::move(new_mask), cudf::UNKNOWN_NULL_COUNT, stream, mr); - } + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, mr, stream); - if (lhs.size() == 0) { return out; } + if (lhs.is_empty()) return out; auto out_view = out->mutable_view(); binops::jit::binary_operation(out_view, lhs, rhs, op, stream); @@ -320,30 +689,22 @@ std::unique_ptr binary_operation(column_view const& lhs, rmm::mr::device_memory_resource* mr, cudaStream_t stream) { - CUDF_EXPECTS((lhs.size() == rhs.size()), "Column sizes don't match"); + CUDF_EXPECTS(lhs.size() == rhs.size(), "Column sizes don't match"); - if ((lhs.type().id() == type_id::STRING) && (rhs.type().id() == type_id::STRING)) { + if (lhs.type().id() == type_id::STRING and rhs.type().id() == type_id::STRING) return binops::compiled::binary_operation(lhs, rhs, op, output_type, mr, stream); - } + + if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) + return fixed_point_binary_operation(lhs, rhs, op, mr, stream); // Check for datatype CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype"); - 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)) { - 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); - } - }(); + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, mr, stream); - // Check for 0 sized data - if (lhs.size() == 0 || rhs.size() == 0) { return out; } + if (lhs.is_empty() or rhs.is_empty()) return out; auto out_view = out->mutable_view(); binops::jit::binary_operation(out_view, lhs, rhs, op, stream); @@ -359,7 +720,8 @@ std::unique_ptr binary_operation(column_view const& lhs, { // Check for datatype auto is_type_supported_ptx = [](data_type type) -> bool { - return is_fixed_width(type) and type.id() != type_id::INT8; // Numba PTX doesn't support int8 + return is_fixed_width(type) and not is_fixed_point(type) and + type.id() != type_id::INT8; // Numba PTX doesn't support int8 }; CUDF_EXPECTS(is_type_supported_ptx(lhs.type()), "Invalid/Unsupported lhs datatype"); @@ -373,7 +735,7 @@ std::unique_ptr binary_operation(column_view const& lhs, 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; } + if (lhs.is_empty() or rhs.is_empty()) return out; auto out_view = out->mutable_view(); binops::jit::binary_operation(out_view, lhs, rhs, ptx, stream); diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu index 945947cbc64..a77f16b30ce 100644 --- a/cpp/src/binaryop/compiled/binary_ops.cu +++ b/cpp/src/binaryop/compiled/binary_ops.cu @@ -423,7 +423,7 @@ std::unique_ptr binary_operation(scalar const& lhs, // hard-coded to only work with cudf::string_view so we don't explode compile times CUDF_EXPECTS(lhs.type().id() == cudf::type_id::STRING, "Invalid/Unsupported lhs datatype"); CUDF_EXPECTS(rhs.type().id() == cudf::type_id::STRING, "Invalid/Unsupported rhs datatype"); - if (null_using_binop(op)) { + if (is_null_dependent(op)) { if (rhs.size() == 0) return cudf::make_empty_column(output_type); auto rhs_device_view = cudf::column_device_view::create(rhs, stream); return null_considering_binop{}(lhs, *rhs_device_view, op, output_type, rhs.size(), mr, stream); @@ -445,7 +445,7 @@ std::unique_ptr binary_operation(column_view const& lhs, // hard-coded to only work with cudf::string_view so we don't explode compile times CUDF_EXPECTS(lhs.type().id() == cudf::type_id::STRING, "Invalid/Unsupported lhs datatype"); CUDF_EXPECTS(rhs.type().id() == cudf::type_id::STRING, "Invalid/Unsupported rhs datatype"); - if (null_using_binop(op)) { + if (is_null_dependent(op)) { if (lhs.size() == 0) return cudf::make_empty_column(output_type); auto lhs_device_view = cudf::column_device_view::create(lhs, stream); return null_considering_binop{}(*lhs_device_view, rhs, op, output_type, lhs.size(), mr, stream); @@ -466,7 +466,7 @@ std::unique_ptr binary_operation(column_view const& lhs, // hard-coded to only work with cudf::string_view so we don't explode compile times CUDF_EXPECTS(lhs.type().id() == cudf::type_id::STRING, "Invalid/Unsupported lhs datatype"); CUDF_EXPECTS(rhs.type().id() == cudf::type_id::STRING, "Invalid/Unsupported rhs datatype"); - if (null_using_binop(op)) { + if (is_null_dependent(op)) { CUDF_EXPECTS(lhs.size() == rhs.size(), "Column sizes do not match"); if (lhs.size() == 0) return cudf::make_empty_column(output_type); auto lhs_device_view = cudf::column_device_view::create(lhs, stream); diff --git a/cpp/src/binaryop/compiled/binary_ops.hpp b/cpp/src/binaryop/compiled/binary_ops.hpp index b0079b076ef..3e6203ce8dd 100644 --- a/cpp/src/binaryop/compiled/binary_ops.hpp +++ b/cpp/src/binaryop/compiled/binary_ops.hpp @@ -35,7 +35,7 @@ rmm::device_buffer scalar_col_valid_mask_and(column_view const& col, * @brief Does the binop need to know if an operand is null/invalid to perform special * processing? */ -inline bool null_using_binop(binary_operator op) +inline bool is_null_dependent(binary_operator op) { return op == binary_operator::NULL_EQUALS || op == binary_operator::NULL_MIN || op == binary_operator::NULL_MAX; diff --git a/cpp/src/column/column_factories.cpp b/cpp/src/column/column_factories.cpp index b3dafbe574a..648e1a14708 100644 --- a/cpp/src/column/column_factories.cpp +++ b/cpp/src/column/column_factories.cpp @@ -18,6 +18,7 @@ #include #include #include +#include #include #include #include @@ -30,24 +31,33 @@ namespace cudf { namespace { struct size_of_helper { - template - constexpr std::enable_if_t(), int> operator()() const + cudf::data_type type; + template ()>* = nullptr> + constexpr int operator()() const { CUDF_FAIL("Invalid, non fixed-width element type."); } - template - constexpr std::enable_if_t(), int> operator()() const noexcept + template () && not is_fixed_point()>* = nullptr> + constexpr int operator()() const noexcept { return sizeof(T); } + + template ()>* = nullptr> + constexpr int operator()() const noexcept + { + // Only want the sizeof fixed_point::Rep as fixed_point::scale is stored in data_type + return sizeof(typename T::rep); + } }; } // namespace std::size_t size_of(data_type element_type) { CUDF_EXPECTS(is_fixed_width(element_type), "Invalid element type."); - return cudf::type_dispatcher(element_type, size_of_helper{}); + return cudf::type_dispatcher(element_type, size_of_helper{element_type}); } // Empty column of specified type @@ -76,6 +86,24 @@ std::unique_ptr make_numeric_column(data_type type, std::vector>{}); } +// Allocate storage for a specified number of numeric elements +std::unique_ptr make_fixed_point_column(data_type type, + size_type size, + mask_state state, + cudaStream_t stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + CUDF_EXPECTS(is_fixed_point(type), "Invalid, non-fixed_point type."); + + return std::make_unique(type, + size, + rmm::device_buffer{size * cudf::size_of(type), stream, mr}, + create_null_mask(size, state, stream, mr), + state_null_count(state, size), + std::vector>{}); +} + // Allocate storage for a specified number of timestamp elements std::unique_ptr make_timestamp_column(data_type type, size_type size, @@ -122,12 +150,12 @@ std::unique_ptr make_fixed_width_column(data_type type, CUDF_FUNC_RANGE(); CUDF_EXPECTS(is_fixed_width(type), "Invalid, non-fixed-width type."); - if (is_timestamp(type)) { - return make_timestamp_column(type, size, state, stream, mr); - } else if (is_duration(type)) { - return make_duration_column(type, size, state, stream, mr); - } - return make_numeric_column(type, size, state, stream, mr); + // clang-format off + if (is_timestamp (type)) return make_timestamp_column (type, size, state, stream, mr); + else if (is_duration (type)) return make_duration_column (type, size, state, stream, mr); + else if (is_fixed_point(type)) return make_fixed_point_column(type, size, state, stream, mr); + else return make_numeric_column (type, size, state, stream, mr); + /// clang-format on } struct column_from_scalar_dispatch { diff --git a/cpp/src/copying/concatenate.cu b/cpp/src/copying/concatenate.cu index 6c475c32e58..225e08eb1a8 100644 --- a/cpp/src/copying/concatenate.cu +++ b/cpp/src/copying/concatenate.cu @@ -294,11 +294,13 @@ struct concatenate_dispatch { bool const has_nulls = std::any_of(views.cbegin(), views.cend(), [](auto const& col) { return col.has_nulls(); }); + using Type = device_storage_type_t; + // Use a heuristic to guess when the fused kernel will be faster if (use_fused_kernel_heuristic(has_nulls, views.size())) { - return fused_concatenate(views, has_nulls, mr, stream); + return fused_concatenate(views, has_nulls, mr, stream); } else { - return for_each_concatenate(views, has_nulls, mr, stream); + return for_each_concatenate(views, has_nulls, mr, stream); } } }; diff --git a/cpp/src/copying/copy_range.cu b/cpp/src/copying/copy_range.cu index 2f6df0482bb..812867ba3ca 100644 --- a/cpp/src/copying/copy_range.cu +++ b/cpp/src/copying/copy_range.cu @@ -148,6 +148,28 @@ std::unique_ptr out_of_place_copy_range_dispatch::operator() +std::unique_ptr out_of_place_copy_range_dispatch::operator()( + cudf::size_type source_begin, + cudf::size_type source_end, + cudf::size_type target_begin, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) +{ + CUDF_FAIL("decimal64 type not supported"); +} + +template <> +std::unique_ptr out_of_place_copy_range_dispatch::operator()( + cudf::size_type source_begin, + cudf::size_type source_end, + cudf::size_type target_begin, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) +{ + CUDF_FAIL("decimal32 type not supported"); +} + template <> std::unique_ptr out_of_place_copy_range_dispatch::operator()( cudf::size_type source_begin, diff --git a/cpp/src/copying/scatter.cu b/cpp/src/copying/scatter.cu index cb79bf3c0d1..6244800a51a 100644 --- a/cpp/src/copying/scatter.cu +++ b/cpp/src/copying/scatter.cu @@ -135,8 +135,10 @@ struct column_scalar_scatterer_impl { auto result = std::make_unique(target, stream, mr); auto result_view = result->mutable_view(); + using Type = device_storage_type_t; + // Use permutation iterator with constant index to dereference scalar data - auto scalar_impl = static_cast*>(source.get()); + auto scalar_impl = static_cast*>(source.get()); auto scalar_iter = thrust::make_permutation_iterator(scalar_impl->data(), thrust::make_constant_iterator(0)); @@ -144,7 +146,7 @@ struct column_scalar_scatterer_impl { scalar_iter, scalar_iter + scatter_rows, scatter_iter, - result_view.begin()); + result_view.begin()); return result; } diff --git a/cpp/src/copying/shift.cu b/cpp/src/copying/shift.cu index 8d70c36c7ba..1beda950149 100644 --- a/cpp/src/copying/shift.cu +++ b/cpp/src/copying/shift.cu @@ -56,7 +56,8 @@ struct shift_functor { rmm::mr::device_memory_resource* mr, cudaStream_t stream) { - using ScalarType = cudf::scalar_type_t; + using Type = device_storage_type_t; + using ScalarType = cudf::scalar_type_t; auto& scalar = static_cast(fill_value); auto device_input = column_device_view::create(input); @@ -83,7 +84,7 @@ struct shift_functor { output->set_null_count(std::get<1>(mask_pair)); } - auto data = device_output->data(); + auto data = device_output->data(); // avoid assigning elements we know to be invalid. if (not scalar.is_valid()) { @@ -98,7 +99,7 @@ struct shift_functor { auto func_value = [size, offset, fill = scalar.data(), input = *device_input] __device__(size_type idx) { auto src_idx = idx - offset; - return out_of_bounds(size, src_idx) ? *fill : input.element(src_idx); + return out_of_bounds(size, src_idx) ? *fill : input.element(src_idx); }; thrust::transform( diff --git a/cpp/src/copying/split.cpp b/cpp/src/copying/split.cpp index 052275870ac..97520800408 100644 --- a/cpp/src/copying/split.cpp +++ b/cpp/src/copying/split.cpp @@ -26,7 +26,7 @@ namespace { template std::vector split(T const& input, size_type column_size, std::vector const& splits) { - if (splits.size() == 0 or column_size == 0) { return std::vector{input}; } + if (splits.empty() or column_size == 0) { return std::vector{input}; } CUDF_EXPECTS(splits.back() <= column_size, "splits can't exceed size of input columns"); // If the size is not zero, the split will always start at `0` diff --git a/cpp/src/dictionary/search.cu b/cpp/src/dictionary/search.cu index cd2258da2db..a63c582a86f 100644 --- a/cpp/src/dictionary/search.cu +++ b/cpp/src/dictionary/search.cu @@ -72,17 +72,20 @@ struct find_index_fn { return type_dispatcher(input.indices().type(), dispatch_scalar_index{}, 0, false, stream, mr); CUDF_EXPECTS(input.keys().type() == key.type(), "search key type must match dictionary keys type"); - auto keys_view = column_device_view::create(input.keys(), stream); - auto find_key = static_cast const&>(key).value(stream); + + using Type = device_storage_type_t; + using ScalarType = cudf::scalar_type_t; + auto find_key = static_cast(key).value(stream); + auto keys_view = column_device_view::create(input.keys(), stream); auto iter = thrust::equal_range(thrust::device, // segfaults: rmm::exec_policy(stream)->on(stream) and // thrust::cuda::par.on(stream) - keys_view->begin(), - keys_view->end(), + keys_view->begin(), + keys_view->end(), find_key); return type_dispatcher(input.indices().type(), dispatch_scalar_index{}, - thrust::distance(keys_view->begin(), iter.first), + thrust::distance(keys_view->begin(), iter.first), (thrust::distance(iter.first, iter.second) > 0), stream, mr); @@ -140,15 +143,18 @@ struct find_insert_index_fn { return type_dispatcher(input.indices().type(), dispatch_scalar_index{}, 0, false, stream, mr); CUDF_EXPECTS(input.keys().type() == key.type(), "search key type must match dictionary keys type"); - auto keys_view = column_device_view::create(input.keys(), stream); - auto find_key = static_cast const&>(key).value(stream); - auto iter = thrust::lower_bound(rmm::exec_policy(stream)->on(stream), - keys_view->begin(), - keys_view->end(), + + using Type = device_storage_type_t; + using ScalarType = cudf::scalar_type_t; + auto find_key = static_cast(key).value(stream); + auto keys_view = column_device_view::create(input.keys(), stream); + auto iter = thrust::lower_bound(rmm::exec_policy(stream)->on(stream), + keys_view->begin(), + keys_view->end(), find_key); return type_dispatcher(input.indices().type(), dispatch_scalar_index{}, - thrust::distance(keys_view->begin(), iter), + thrust::distance(keys_view->begin(), iter), true, stream, mr); diff --git a/cpp/src/io/avro/avro.cpp b/cpp/src/io/avro/avro.cpp index 5260d0b8113..fd3b4095e6f 100644 --- a/cpp/src/io/avro/avro.cpp +++ b/cpp/src/io/avro/avro.cpp @@ -317,7 +317,7 @@ bool schema_parser::parse(std::vector &schema, const std::string & /** * @Brief Parse a string * - * @returns return parsed string, consuming the terminating quote + * @returns parsed string, consuming the terminating quote */ std::string schema_parser::get_str() { diff --git a/cpp/src/io/utilities/column_buffer.hpp b/cpp/src/io/utilities/column_buffer.hpp index ac41ac1bcf7..90bdc42804c 100644 --- a/cpp/src/io/utilities/column_buffer.hpp +++ b/cpp/src/io/utilities/column_buffer.hpp @@ -61,8 +61,8 @@ inline rmm::device_buffer create_data( } /** - * @brief Class for holding device memory buffers to column data that will be - * eventually used create to create a column. + * @brief Class for holding device memory buffers to column data that eventually + * will be used to create a column. */ struct column_buffer { // there is a potential bug here. In the decoding step, the buffer of diff --git a/cpp/src/jit/type.cpp b/cpp/src/jit/type.cpp index 60ebc97e20d..d71e2eb4df8 100644 --- a/cpp/src/jit/type.cpp +++ b/cpp/src/jit/type.cpp @@ -70,9 +70,9 @@ std::string get_type_name(data_type type) { // TODO: Remove in JIT type utils PR switch (type.id()) { - case type_id::BOOL8: return CUDF_STRINGIFY(bool); - case type_id::LIST: return CUDF_STRINGIFY(List); + case type_id::DECIMAL32: return CUDF_STRINGIFY(int32_t); + case type_id::DECIMAL64: return CUDF_STRINGIFY(int64_t); default: break; } diff --git a/cpp/src/merge/merge.cu b/cpp/src/merge/merge.cu index 2e346c0f9c0..90250278203 100644 --- a/cpp/src/merge/merge.cu +++ b/cpp/src/merge/merge.cu @@ -236,7 +236,7 @@ struct column_merger { index_vector const& row_order, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource(), cudaStream_t stream = nullptr) - : dv_row_order_(row_order), mr_(mr), stream_(stream) + : row_order_(row_order), mr_(mr), stream_(stream) { } @@ -248,18 +248,14 @@ struct column_merger { auto lsz = lcol.size(); auto merged_size = lsz + rcol.size(); auto type = lcol.type(); + auto merged_col = lcol.has_nulls() ? cudf::allocate_like(lcol, merged_size) + : cudf::allocate_like(rcol, merged_size); - std::unique_ptr p_merged_col{nullptr}; - if (lcol.has_nulls()) - p_merged_col = cudf::allocate_like(lcol, merged_size); - else - p_merged_col = cudf::allocate_like(rcol, merged_size); - - //"gather" data from lcol, rcol according to dv_row_order_ "map" + //"gather" data from lcol, rcol according to row_order_ "map" //(directly calling gather() won't work because // lcol, rcol indices overlap!) // - cudf::mutable_column_view merged_view = p_merged_col->mutable_view(); + cudf::mutable_column_view merged_view = merged_col->mutable_view(); // initialize null_mask to all valid: // @@ -272,12 +268,14 @@ struct column_merger { // set the null count: // - p_merged_col->set_null_count(lcol.null_count() + rcol.null_count()); + merged_col->set_null_count(lcol.null_count() + rcol.null_count()); + + using Type = device_storage_type_t; // to resolve view.data()'s types use: Element // - Element const* p_d_lcol = lcol.data(); - Element const* p_d_rcol = rcol.data(); + auto const d_lcol = lcol.data(); + auto const d_rcol = rcol.data(); auto exe_pol = rmm::exec_policy(stream_); @@ -286,15 +284,14 @@ struct column_merger { // from lcol or rcol, depending on side; // thrust::transform(exe_pol->on(stream_), - dv_row_order_.begin(), - dv_row_order_.end(), - merged_view.begin(), - [p_d_lcol, p_d_rcol] __device__(index_type const& index_pair) { + row_order_.begin(), + row_order_.end(), + merged_view.begin(), + [d_lcol, d_rcol] __device__(index_type const& index_pair) { + // When C++17, use structure bindings auto side = thrust::get<0>(index_pair); auto index = thrust::get<1>(index_pair); - - Element val = (side == side::LEFT ? p_d_lcol[index] : p_d_rcol[index]); - return val; + return side == side::LEFT ? d_lcol[index] : d_rcol[index]; }); // CAVEAT: conditional call below is erroneous without @@ -303,14 +300,14 @@ struct column_merger { if (lcol.has_nulls() || rcol.has_nulls()) { // resolve null mask: // - materialize_bitmask(lcol, rcol, merged_view, dv_row_order_.data().get(), stream_); + materialize_bitmask(lcol, rcol, merged_view, row_order_.data().get(), stream_); } - return p_merged_col; + return merged_col; } private: - index_vector const& dv_row_order_; + index_vector const& row_order_; rmm::mr::device_memory_resource* mr_; cudaStream_t stream_; }; @@ -322,13 +319,13 @@ std::unique_ptr column_merger::operator()(column_view { auto column = strings::detail::merge(strings_column_view(lcol), strings_column_view(rcol), - dv_row_order_.begin(), - dv_row_order_.end(), + row_order_.begin(), + row_order_.end(), mr_, stream_); if (lcol.has_nulls() || rcol.has_nulls()) { auto merged_view = column->mutable_view(); - materialize_bitmask(lcol, rcol, merged_view, dv_row_order_.data().get(), stream_); + materialize_bitmask(lcol, rcol, merged_view, row_order_.data().get(), stream_); } return column; } diff --git a/cpp/src/reductions/simple.cuh b/cpp/src/reductions/simple.cuh index f69c38bb6cb..980b709e241 100644 --- a/cpp/src/reductions/simple.cuh +++ b/cpp/src/reductions/simple.cuh @@ -81,8 +81,8 @@ struct result_type_dispatcher { return cudf::is_convertible::value && (std::is_arithmetic::value || std::is_same::value || - std::is_same::value || - cudf::is_fixed_point()) && + std::is_same::value) && + !cudf::is_fixed_point() && !std::is_same::value && !std::is_same::value; } diff --git a/cpp/src/replace/clamp.cu b/cpp/src/replace/clamp.cu index 731491430fa..38af16ed5e2 100644 --- a/cpp/src/replace/clamp.cu +++ b/cpp/src/replace/clamp.cu @@ -290,7 +290,7 @@ std::unique_ptr dispatch_clamp::operator()( rmm::mr::device_memory_resource* mr, cudaStream_t stream) { - CUDF_FAIL("clamp for decimal32 not supported"); + CUDF_FAIL("clamp for decimal64 not supported"); } template <> diff --git a/cpp/src/replace/nulls.cu b/cpp/src/replace/nulls.cu index 4ac27e441d8..f38fcc5904d 100644 --- a/cpp/src/replace/nulls.cu +++ b/cpp/src/replace/nulls.cu @@ -285,7 +285,8 @@ struct replace_nulls_functor { * `replace_nulls` with the appropriate data types. */ struct replace_nulls_scalar_kernel_forwarder { - template ()>* = nullptr> + template ()>* = nullptr> std::unique_ptr operator()(cudf::column_view const& input, cudf::scalar const& replacement, rmm::mr::device_memory_resource* mr, @@ -296,16 +297,17 @@ struct replace_nulls_scalar_kernel_forwarder { cudf::allocate_like(input, cudf::mask_allocation_policy::NEVER, mr); auto output_view = output->mutable_view(); + using Type = cudf::device_storage_type_t; using ScalarType = cudf::scalar_type_t; auto s1 = static_cast(replacement); auto device_in = cudf::column_device_view::create(input); - replace_nulls_functor func(s1.data()); + auto func = replace_nulls_functor{s1.data()}; thrust::transform(rmm::exec_policy(stream)->on(stream), - input.data(), - input.data() + input.size(), + input.data(), + input.data() + input.size(), cudf::detail::make_validity_iterator(*device_in), - output_view.data(), + output_view.data(), func); return output; } diff --git a/cpp/src/replace/replace.cu b/cpp/src/replace/replace.cu index 9b287ac003f..07a9f0fab9f 100644 --- a/cpp/src/replace/replace.cu +++ b/cpp/src/replace/replace.cu @@ -240,7 +240,9 @@ __global__ void replace_kernel(cudf::column_device_view input, cudf::column_device_view values_to_replace, cudf::column_device_view replacement) { - T* __restrict__ output_data = output.data(); + using Type = cudf::device_storage_type_t; + + Type* __restrict__ output_data = output.data(); cudf::size_type i = blockIdx.x * blockDim.x + threadIdx.x; @@ -257,12 +259,12 @@ __global__ void replace_kernel(cudf::column_device_view input, output_is_valid = input_is_valid; } if (input_is_valid) - thrust::tie(output_data[i], output_is_valid) = get_new_value( + thrust::tie(output_data[i], output_is_valid) = get_new_value( i, - input.data(), - values_to_replace.data(), - values_to_replace.data() + values_to_replace.size(), - replacement.data(), + input.data(), + values_to_replace.data(), + values_to_replace.data() + values_to_replace.size(), + replacement.data(), replacement.null_mask()); /* output valid counts calculations*/ @@ -301,47 +303,39 @@ struct replace_kernel_forwarder { rmm::device_scalar valid_counter(0, stream); cudf::size_type* valid_count = valid_counter.data(); - auto replace = replace_kernel; - if (input_col.has_nulls()) { - if (replacement_values.has_nulls()) { - replace = replace_kernel; - } else { - replace = replace_kernel; - } - } else { - if (replacement_values.has_nulls()) { - replace = replace_kernel; - } else { - replace = replace_kernel; - } - } - - std::unique_ptr output; - if (input_col.has_nulls() || replacement_values.has_nulls()) { - output = cudf::detail::allocate_like( - input_col, input_col.size(), cudf::mask_allocation_policy::ALWAYS, mr, stream); - } else { - output = cudf::detail::allocate_like( - input_col, input_col.size(), cudf::mask_allocation_policy::NEVER, mr, stream); - } - - cudf::mutable_column_view outputView = output->mutable_view(); - - cudf::detail::grid_1d grid{outputView.size(), BLOCK_SIZE, 1}; + auto replace = [&] { + if (input_col.has_nulls()) + return replacement_values.has_nulls() ? replace_kernel + : replace_kernel; + else + return replacement_values.has_nulls() ? replace_kernel + : replace_kernel; + }(); + + auto output = [&] { + auto const mask_allocation_policy = input_col.has_nulls() || replacement_values.has_nulls() + ? cudf::mask_allocation_policy::ALWAYS + : cudf::mask_allocation_policy::NEVER; + return cudf::detail::allocate_like( + input_col, input_col.size(), mask_allocation_policy, mr, stream); + }(); + + auto output_view = output->mutable_view(); + auto grid = cudf::detail::grid_1d{output_view.size(), BLOCK_SIZE, 1}; auto device_in = cudf::column_device_view::create(input_col); - auto device_out = cudf::mutable_column_device_view::create(outputView); + auto device_out = cudf::mutable_column_device_view::create(output_view); auto device_values_to_replace = cudf::column_device_view::create(values_to_replace); auto device_replacement_values = cudf::column_device_view::create(replacement_values); replace<<>>(*device_in, *device_out, valid_count, - outputView.size(), + output_view.size(), *device_values_to_replace, *device_replacement_values); - if (outputView.nullable()) { + if (output_view.nullable()) { output->set_null_count(output->size() - valid_counter.value(stream)); } return output; diff --git a/cpp/src/reshape/interleave_columns.cu b/cpp/src/reshape/interleave_columns.cu index c181e092e58..d6c68d56797 100644 --- a/cpp/src/reshape/interleave_columns.cu +++ b/cpp/src/reshape/interleave_columns.cu @@ -133,16 +133,18 @@ struct interleave_columns_functor { auto index_begin = thrust::make_counting_iterator(0); auto index_end = thrust::make_counting_iterator(output_size); + using Type = device_storage_type_t; + auto func_value = [input = *device_input, divisor = input.num_columns()] __device__(size_type idx) { - return input.column(idx % divisor).element(idx / divisor); + return input.column(idx % divisor).element(idx / divisor); }; if (not create_mask) { thrust::transform(rmm::exec_policy(stream)->on(stream), index_begin, index_end, - device_output->data(), + device_output->begin(), func_value); return output; @@ -156,7 +158,7 @@ struct interleave_columns_functor { thrust::transform_if(rmm::exec_policy(stream)->on(stream), index_begin, index_end, - device_output->data(), + device_output->begin(), func_value, func_validity); @@ -180,13 +182,15 @@ std::unique_ptr interleave_columns(table_view const& input, CUDF_FUNC_RANGE(); CUDF_EXPECTS(input.num_columns() > 0, "input must have at least one column to determine dtype."); - auto dtype = input.column(0).type(); - auto output_needs_mask = false; + auto const dtype = input.column(0).type(); - for (auto& col : input) { - CUDF_EXPECTS(dtype == col.type(), "DTYPE mismatch"); - output_needs_mask |= col.nullable(); - } + CUDF_EXPECTS(std::all_of(std::cbegin(input), + std::cend(input), + [dtype](auto const& col) { return dtype == col.type(); }), + "DTYPE mismatch"); + + auto const output_needs_mask = std::any_of( + std::cbegin(input), std::cend(input), [](auto const& col) { return col.nullable(); }); return type_dispatcher(dtype, detail::interleave_columns_functor{}, input, output_needs_mask, mr); } diff --git a/cpp/src/scalar/scalar_factories.cpp b/cpp/src/scalar/scalar_factories.cpp index a0d7082c905..7bad39af717 100644 --- a/cpp/src/scalar/scalar_factories.cpp +++ b/cpp/src/scalar/scalar_factories.cpp @@ -23,16 +23,30 @@ namespace cudf { namespace { struct scalar_construction_helper { - template > - std::enable_if_t(), std::unique_ptr> operator()( - cudaStream_t stream, rmm::mr::device_memory_resource* mr) const + template , + typename std::enable_if_t() and not is_fixed_point()>* = nullptr> + std::unique_ptr operator()(cudaStream_t stream, rmm::mr::device_memory_resource* mr) const { - auto s = new ScalarType(T{}, false, stream, mr); + using Type = device_storage_type_t; + auto s = new ScalarType(Type{}, false, stream, mr); return std::unique_ptr(s); } - template - std::enable_if_t(), std::unique_ptr> operator()(Args... args) const + template , + typename std::enable_if_t()>* = nullptr> + std::unique_ptr operator()(cudaStream_t stream, rmm::mr::device_memory_resource* mr) const + { + using Type = device_storage_type_t; + auto s = new ScalarType(Type{}, numeric::scale_type{0}, false, stream, mr); + return std::unique_ptr(s); + } + + template ()>* = nullptr> + std::unique_ptr operator()(Args... args) const { CUDF_FAIL("Invalid type."); } diff --git a/cpp/src/search/search.cu b/cpp/src/search/search.cu index de82cae40ef..2cdda841e3e 100644 --- a/cpp/src/search/search.cu +++ b/cpp/src/search/search.cu @@ -143,24 +143,25 @@ struct contains_scalar_dispatch { { CUDF_EXPECTS(col.type() == value.type(), "scalar and column types must match"); + using Type = device_storage_type_t; using ScalarType = cudf::scalar_type_t; auto d_col = column_device_view::create(col, stream); auto s = static_cast(&value); if (col.has_nulls()) { auto found_iter = thrust::find(rmm::exec_policy(stream)->on(stream), - d_col->pair_begin(), - d_col->pair_end(), + d_col->pair_begin(), + d_col->pair_end(), thrust::make_pair(s->value(), true)); - return found_iter != d_col->pair_end(); + return found_iter != d_col->pair_end(); } else { - auto found_iter = thrust::find(rmm::exec_policy(stream)->on(stream), - d_col->begin(), - d_col->end(), + auto found_iter = thrust::find(rmm::exec_policy(stream)->on(stream), // + d_col->begin(), + d_col->end(), s->value()); - return found_iter != d_col->end(); + return found_iter != d_col->end(); } } }; diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 7de7c4f735b..500b817bc72 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -22,9 +22,11 @@ #include #include +#include #include #include #include +#include #include namespace cudf { @@ -2047,140 +2049,203 @@ 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 +// struct FixedPointTestBothReps : public cudf::test::BaseFixture { +// }; -template -using wrapper = cudf::test::fixed_width_column_wrapper; -TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); +// template +// using wrapper = cudf::test::fixed_width_column_wrapper; +// TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd) -{ - using namespace numeric; - using decimalXX = TypeParam; +// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd) +// { +// using namespace numeric; +// using decimalXX = TypeParam; - auto const sz = std::size_t{1000}; +// 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); +// auto vec1 = std::vector(sz); +// auto const vec2 = std::vector(sz, decimalXX{2, scale_type{0}}); +// auto expected = std::vector(sz); - std::iota(std::begin(vec1), std::end(vec1), decimalXX{}); +// std::iota(std::begin(vec1), std::end(vec1), decimalXX{1, scale_type{0}}); - std::transform(std::cbegin(vec1), - std::cend(vec1), - std::cbegin(vec2), - std::begin(expected), - std::plus()); +// 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 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()); +// auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_col, result->view()); -} +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_col, result->view()); +// } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiply) -{ - using namespace numeric; - using decimalXX = TypeParam; +// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiply) +// { +// using namespace numeric; +// using decimalXX = TypeParam; - auto const sz = std::size_t{1000}; +// 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); +// auto vec1 = std::vector(sz); +// auto const vec2 = std::vector(sz, decimalXX{2, scale_type{0}}); +// auto expected = std::vector(sz); - std::iota(std::begin(vec1), std::end(vec1), decimalXX{}); +// std::iota(std::begin(vec1), std::end(vec1), decimalXX{1, scale_type{0}}); - std::transform(std::cbegin(vec1), - std::cend(vec1), - std::cbegin(vec2), - std::begin(expected), - std::multiplies()); +// 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 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()); +// auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::MUL, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_col, result->view()); -} +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_col, result->view()); +// } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimple) -{ - using namespace numeric; - using decimalXX = TypeParam; +// template +// using fp_wrapper = cudf::test::fixed_point_column_wrapper; - 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}}; +// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiply2) +// { +// using namespace numeric; +// using decimalXX = TypeParam; +// using RepType = device_storage_type_t; - 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 lhs = fp_wrapper{{11, 22, 33, 44, 55}, scale_type{-1}}; +// auto const rhs = fp_wrapper{{10, 10, 10, 10, 10}, scale_type{0}}; +// auto const expected = fp_wrapper{{110, 220, 330, 440, 550}, scale_type{-1}}; - 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 result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::MUL, {}); - auto const col1 = wrapper(vec1.begin(), vec1.end()); - auto const col2 = wrapper(vec2.begin(), vec2.end()); - auto const expected = wrapper(trues.begin(), trues.end()); +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +// } - auto const result = cudf::binary_operation( - col1, col2, cudf::binary_operator::EQUAL, cudf::data_type{type_id::BOOL8}); +// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpDiv) +// { +// using namespace numeric; +// using decimalXX = TypeParam; +// using RepType = device_storage_type_t; - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); -} +// auto const lhs = fp_wrapper{{10, 30, 50, 70}, scale_type{-1}}; +// auto const rhs = fp_wrapper{{4, 4, 4, 4}, scale_type{0}}; +// auto const expected = fp_wrapper{{3, 8, 13, 18}, scale_type{-1}}; -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualLessGreater) -{ - using namespace numeric; - using decimalXX = TypeParam; +// auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, {}); - auto const sz = std::size_t{1000}; +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +// } - // TESTING binary op ADD +// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd2) +// { +// using namespace numeric; +// using decimalXX = TypeParam; +// using RepType = device_storage_type_t; - auto vec1 = std::vector(sz, decimalXX{0, scale_type{-3}}); - auto vec2 = std::vector(sz, decimalXX{0, scale_type{-1}}); +// auto const lhs = fp_wrapper{{11, 22, 33, 44, 55}, scale_type{-1}}; +// auto const rhs = fp_wrapper{{100, 200, 300, 400, 500}, scale_type{-2}}; +// auto const expected = fp_wrapper{{210, 420, 630, 840, 1050}, scale_type{-2}}; - std::iota(std::begin(vec1), std::end(vec1), decimalXX{1, scale_type{-3}}); +// auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD, {}); - auto const iota_1 = wrapper(vec1.begin(), vec1.end()); - auto const zeros_3 = wrapper(vec2.begin(), vec2.end()); +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +// } - auto const iota_3 = cudf::binary_operation( - zeros_3, iota_1, cudf::binary_operator::ADD, static_cast(zeros_3).type()); +// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd3) +// { +// using namespace numeric; +// using decimalXX = TypeParam; +// using RepType = device_storage_type_t; - CUDF_TEST_EXPECT_COLUMNS_EQUAL(iota_1, iota_3->view()); +// auto const lhs = fp_wrapper{{1100, 2200, 3300, 4400, 5500}, scale_type{-3}}; +// auto const rhs = fp_wrapper{{100, 200, 300, 400, 500}, scale_type{-2}}; +// auto const expected = fp_wrapper{{2100, 4200, 6300, 8400, 10500}, scale_type{-3}}; - // TESTING binary op EQUAL, LESS, GREATER +// auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD, {}); - auto const trues = std::vector(sz, true); - auto const true_col = wrapper(trues.begin(), trues.end()); +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +// } - 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()); +// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiplyScalar) +// { +// using namespace numeric; +// using decimalXX = TypeParam; +// using RepType = device_storage_type_t; - 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 lhs = fp_wrapper{{11, 22, 33, 44, 55}, scale_type{-1}}; +// auto const rhs = make_fixed_point_scalar(100, scale_type{-1}); +// auto const expected = fp_wrapper{{1100, 2200, 3300, 4400, 5500}, scale_type{-2}}; - 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()); -} +// auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::MUL, {}); + +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +// } + +// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimple) +// { +// using namespace numeric; +// using decimalXX = TypeParam; +// using RepType = device_storage_type_t; + +// auto const trues = std::vector(4, true); +// auto const col1 = fp_wrapper{{1, 2, 3, 4}, scale_type{0}}; +// auto const col2 = fp_wrapper{{100, 200, 300, 400}, scale_type{-2}}; +// auto const expected = wrapper(trues.begin(), trues.end()); + +// auto const result = cudf::binary_operation(col1, col2, binary_operator::EQUAL, {}); + +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +// } + +// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualLessGreater) +// { +// using namespace numeric; +// using decimalXX = TypeParam; +// using RepType = device_storage_type_t; + +// auto const sz = std::size_t{1000}; + +// // TESTING binary op ADD + +// auto begin = make_counting_transform_iterator(1, [](auto e) { return e * 1000; }); +// auto const vec1 = std::vector(begin, begin + sz); +// auto const vec2 = std::vector(sz, 0); + +// auto const iota_3 = fp_wrapper(vec1.begin(), vec1.end(), scale_type{-3}); +// auto const zeros_3 = fp_wrapper(vec2.begin(), vec2.end(), scale_type{-1}); + +// auto const iota_3_after_add = cudf::binary_operation(zeros_3, iota_3, binary_operator::ADD, +// {}); + +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(iota_3, iota_3_after_add->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_3, iota_3_after_add->view(), binary_operator::EQUAL, {}); +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(true_col, equal_result->view()); + +// auto const less_result = +// cudf::binary_operation(zeros_3, iota_3_after_add->view(), binary_operator::LESS, {}); +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(true_col, less_result->view()); + +// auto const greater_result = +// cudf::binary_operation(iota_3_after_add->view(), zeros_3, binary_operator::GREATER, {}); +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(true_col, greater_result->view()); +// } } // namespace binop } // namespace test diff --git a/cpp/tests/column/factories_test.cpp b/cpp/tests/column/factories_test.cpp index 6e0bca19178..afd4d6668dd 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::FixedWidthTypesWithoutFixedPoint); +TYPED_TEST_CASE(FixedWidthFactoryTest, cudf::test::FixedWidthTypes); TYPED_TEST(FixedWidthFactoryTest, EmptyNoMask) { diff --git a/cpp/tests/copying/copy_range_tests.cpp b/cpp/tests/copying/copy_range_tests.cpp index 8ab076c8990..fd0c8e455d7 100644 --- a/cpp/tests/copying/copy_range_tests.cpp +++ b/cpp/tests/copying/copy_range_tests.cpp @@ -58,7 +58,7 @@ class CopyRangeTypedTestFixture : public cudf::test::BaseFixture { } }; -TYPED_TEST_CASE(CopyRangeTypedTestFixture, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(CopyRangeTypedTestFixture, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(CopyRangeTypedTestFixture, CopyWithNulls) { diff --git a/cpp/tests/filling/repeat_tests.cpp b/cpp/tests/filling/repeat_tests.cpp index 107659b6bfb..625f3269579 100644 --- a/cpp/tests/filling/repeat_tests.cpp +++ b/cpp/tests/filling/repeat_tests.cpp @@ -52,21 +52,21 @@ TYPED_TEST(RepeatTypedTestFixture, RepeatScalarCount) constexpr cudf::size_type num_values{10}; constexpr cudf::size_type repeat_count{10}; - cudf::test::fixed_width_column_wrapper input( + auto input = cudf::test::fixed_width_column_wrapper( thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + num_values); static_assert(repeat_count > 0, "repeat_count should be larger than 0."); auto expected_elements = cudf::test::make_counting_transform_iterator( 0, [repeat_count](auto i) { return i / repeat_count; }); - cudf::test::fixed_width_column_wrapper - expected(expected_elements, expected_elements + num_values * repeat_count); + auto expected = + cudf::test::fixed_width_column_wrapper( + expected_elements, expected_elements + num_values * repeat_count); - cudf::table_view input_table{{input}}; - - auto p_ret = cudf::repeat(input_table, repeat_count); + auto input_table = cudf::table_view{{input}}; + auto const p_ret = cudf::repeat(input_table, repeat_count); EXPECT_EQ(p_ret->num_columns(), 1); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(p_ret->view().column(0), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(p_ret->view().column(0), expected, true); } TYPED_TEST(RepeatTypedTestFixture, RepeatColumnCount) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index dc8bee438e6..a1ccab22356 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -15,10 +15,15 @@ */ #include +#include #include #include +#include #include +#include +#include +#include #include #include @@ -536,4 +541,43 @@ TEST_F(FixedPointTest, DecimalXXThrustOnDevice) EXPECT_EQ(vec2, vec3); } +TYPED_TEST(FixedPointTestBothReps, FixedPointColumnWrapper) +{ + using namespace numeric; + using decimalXX = fixed_point; + using RepType = TypeParam; + + // fixed_point_column_wrapper + auto const w = cudf::test::fixed_point_column_wrapper{{1, 2, 3, 4}, scale_type{0}}; + + // fixed_width_column_wrapper + 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 vec = std::vector{ONE, TWO, THREE, FOUR}; + auto const col = cudf::test::fixed_width_column_wrapper(vec.begin(), vec.end()); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(col, w); +} + +TYPED_TEST(FixedPointTestBothReps, NoScaleOrWrongTypeID) +{ + auto const null_mask = cudf::create_null_mask(0, cudf::mask_state::ALL_NULL); + + EXPECT_THROW(cudf::make_fixed_point_column(cudf::data_type{cudf::type_id::INT32}, 0, null_mask), + cudf::logic_error); +} + +TYPED_TEST(FixedPointTestBothReps, SimpleFixedPointColumnWrapper) +{ + using decimalXX = fixed_point; + + auto const a = cudf::test::fixed_point_column_wrapper{{11, 22, 33}, scale_type{-1}}; + auto const b = cudf::test::fixed_point_column_wrapper{{110, 220, 330}, scale_type{-2}}; + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(a, b); +} + CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/replace/replace_tests.cpp b/cpp/tests/replace/replace_tests.cpp index 9c83585b28c..51b0b3b9111 100644 --- a/cpp/tests/replace/replace_tests.cpp +++ b/cpp/tests/replace/replace_tests.cpp @@ -23,11 +23,12 @@ #include #include -#include #include #include +#include "cudf/fixed_point/fixed_point.hpp" #include +#include #include #include @@ -548,10 +549,10 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointReplace) 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 mod2 = [&](auto e) { return e % 2 ? ONE : TWO; }; + auto transform_begin = cudf::test::make_counting_transform_iterator(0, mod2); + auto const vec1 = std::vector(transform_begin, transform_begin + sz); + auto const vec2 = std::vector(sz, TWO); auto const to_replace = std::vector{ONE}; auto const replacement = std::vector{TWO}; diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 964da18cb8b..654df7589e6 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -25,6 +25,7 @@ #include #include +#include "cudf/utilities/traits.hpp" using namespace cudf::test; @@ -166,7 +167,7 @@ TYPED_TEST(InterleaveColumnsTest, MismatchedDtypes) { using T = TypeParam; - if (not std::is_same::value) { + if (not std::is_same::value and not cudf::is_fixed_point()) { fixed_width_column_wrapper input_a({1, 4, 7}, {1, 0, 1}); fixed_width_column_wrapper input_b({2, 5, 8}, {0, 1, 0}); diff --git a/cpp/tests/scalar/factories_test.cpp b/cpp/tests/scalar/factories_test.cpp index bfb69c9965d..aee99e7a5bf 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) { @@ -135,4 +135,26 @@ TYPED_TEST(FixedWidthScalarFactory, ValueProvided) EXPECT_TRUE(s->is_valid()); } +template +struct FixedPointScalarFactory : public ScalarFactoryTest { +}; + +TYPED_TEST_CASE(FixedPointScalarFactory, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointScalarFactory, ValueProvided) +{ + using namespace numeric; + using decimalXX = TypeParam; + + auto const rep_value = static_cast(123); + auto const s = + cudf::make_fixed_point_scalar(123, scale_type{-2}, this->stream(), this->mr()); + auto const fp_s = static_cast*>(s.get()); + + EXPECT_EQ(s->type(), cudf::data_type{cudf::type_to_id()}); + EXPECT_EQ(fp_s->value(), rep_value); + EXPECT_TRUE(fp_s->is_valid()); + EXPECT_TRUE(s->is_valid()); +} + CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/scalar/scalar_test.cpp b/cpp/tests/scalar/scalar_test.cpp index d96b5fcebb7..e392b0273a0 100644 --- a/cpp/tests/scalar/scalar_test.cpp +++ b/cpp/tests/scalar/scalar_test.cpp @@ -29,11 +29,17 @@ template struct TypedScalarTest : public cudf::test::BaseFixture { }; +template +struct TypedScalarTestWithoutFixedPoint : public cudf::test::BaseFixture { +}; + TYPED_TEST_CASE(TypedScalarTest, cudf::test::FixedWidthTypes); +TYPED_TEST_CASE(TypedScalarTestWithoutFixedPoint, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(TypedScalarTest, DefaultValidity) { - TypeParam value = cudf::test::make_type_param_scalar(7); + using Type = cudf::device_storage_type_t; + Type value = cudf::test::make_type_param_scalar(7); cudf::scalar_type_t s(value); EXPECT_TRUE(s.is_valid()); @@ -48,7 +54,7 @@ TYPED_TEST(TypedScalarTest, ConstructNull) EXPECT_FALSE(s.is_valid()); } -TYPED_TEST(TypedScalarTest, SetValue) +TYPED_TEST(TypedScalarTestWithoutFixedPoint, SetValue) { TypeParam value = cudf::test::make_type_param_scalar(9); cudf::scalar_type_t s; @@ -58,7 +64,7 @@ TYPED_TEST(TypedScalarTest, SetValue) EXPECT_EQ(value, s.value()); } -TYPED_TEST(TypedScalarTest, SetNull) +TYPED_TEST(TypedScalarTestWithoutFixedPoint, SetNull) { TypeParam value = cudf::test::make_type_param_scalar(6); cudf::scalar_type_t s; @@ -70,7 +76,8 @@ TYPED_TEST(TypedScalarTest, SetNull) TYPED_TEST(TypedScalarTest, CopyConstructor) { - TypeParam value = cudf::test::make_type_param_scalar(8); + using Type = cudf::device_storage_type_t; + Type value = cudf::test::make_type_param_scalar(8); cudf::scalar_type_t s(value); auto s2 = s; diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index 64e383f74e8..9bfaad7e9d0 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -25,7 +25,8 @@ #include #include #include -#include +#include "cudf/utilities/type_dispatcher.hpp" + #include #include #include @@ -33,13 +34,12 @@ #include -#include - #include #include #include #include +#include namespace cudf { namespace test { @@ -164,13 +164,12 @@ std::string differences_message(thrust::device_vector const& differences, auto diff_table = cudf::gather(source_table, diff_column); // Need to pull back the differences - std::vector h_left_strings = to_strings(diff_table->get_column(0)); - std::vector h_right_strings = to_strings(diff_table->get_column(1)); + auto const h_left_strings = to_strings(diff_table->get_column(0)); + auto const h_right_strings = to_strings(diff_table->get_column(1)); - for (size_t i = 0; i < differences.size(); ++i) { + for (size_t i = 0; i < differences.size(); ++i) buffer << depth_str << "lhs[" << differences[i] << "] = " << h_left_strings[i] << ", rhs[" << differences[i] << "] = " << h_right_strings[i] << std::endl; - } return buffer.str(); } else { @@ -564,15 +563,10 @@ struct column_view_printer { std::string const& indent) { auto const h_data = cudf::test::to_host(col); - - out.resize(col.size()); - std::transform(thrust::make_counting_iterator(size_type{0}), - thrust::make_counting_iterator(col.size()), - out.begin(), - [&](auto idx) { - auto const d = static_cast(h_data.first[idx]); - return std::to_string(d); - }); + std::transform(std::cbegin(h_data.first), + std::cend(h_data.first), + std::back_inserter(out), + [](auto const& fp) { return std::to_string(static_cast(fp)); }); } template struct ColumnUtilitiesTestFloatingPoint : public cudf::test::BaseFixture { }; +template +struct ColumnUtilitiesTestFixedPoint : public cudf::test::BaseFixture { +}; + TYPED_TEST_CASE(ColumnUtilitiesTest, cudf::test::FixedWidthTypes); TYPED_TEST_CASE(ColumnUtilitiesTestIntegral, cudf::test::IntegralTypes); TYPED_TEST_CASE(ColumnUtilitiesTestFloatingPoint, cudf::test::FloatingPointTypes); +TYPED_TEST_CASE(ColumnUtilitiesTestFixedPoint, cudf::test::FixedPointTypes); TYPED_TEST(ColumnUtilitiesTest, NonNullableToHost) { @@ -68,15 +73,15 @@ TYPED_TEST(ColumnUtilitiesTest, NonNullableToHostWithOffset) auto sequence = cudf::test::make_counting_transform_iterator( 0, [](auto i) { return cudf::test::make_type_param_scalar(i); }); - auto size = this->size(); - auto split = 2; + auto const size = this->size(); + auto const split = 2; - std::vector data(sequence, sequence + size); - std::vector expected_data(sequence + split, sequence + size); - cudf::test::fixed_width_column_wrapper col(data.begin(), data.end()); + auto data = std::vector(sequence, sequence + size); + auto expected_data = std::vector(sequence + split, sequence + size); + auto col = cudf::test::fixed_width_column_wrapper(data.begin(), data.end()); - std::vector splits{split}; - std::vector result = cudf::split(col, splits); + auto const splits = std::vector{split}; + auto result = cudf::split(col, splits); auto host_data = cudf::test::to_host(result.back()); @@ -271,4 +276,49 @@ TEST_F(ColumnUtilitiesStringsTest, StringsToString) EXPECT_EQ(cudf::test::to_string(strings, delimiter), tmp.str()); } +TYPED_TEST(ColumnUtilitiesTestFixedPoint, NonNullableToHost) +{ + using namespace numeric; + using decimalXX = TypeParam; + using rep = cudf::device_storage_type_t; + + auto const scale = scale_type{-2}; + auto to_fp = [&](auto i) { return decimalXX{i, scale}; }; + auto to_rep = [](auto i) { return i * 100; }; + auto fps = cudf::test::make_counting_transform_iterator(0, to_fp); + auto reps = cudf::test::make_counting_transform_iterator(0, to_rep); + + auto const size = 1000; + auto const expected = std::vector(fps, fps + size); + auto const col = cudf::test::fixed_point_column_wrapper(reps, reps + size, scale); + auto const host_data = cudf::test::to_host(col); + + EXPECT_TRUE(std::equal(expected.begin(), expected.end(), host_data.first.begin())); +} + +TYPED_TEST(ColumnUtilitiesTestFixedPoint, NonNullableToHostWithOffset) +{ + using namespace numeric; + using decimalXX = TypeParam; + using rep = cudf::device_storage_type_t; + + auto const scale = scale_type{-2}; + auto to_fp = [&](auto i) { return decimalXX{i, scale}; }; + auto to_rep = [](auto i) { return i * 100; }; + auto fps = cudf::test::make_counting_transform_iterator(0, to_fp); + auto reps = cudf::test::make_counting_transform_iterator(0, to_rep); + + auto const size = 1000; + auto const split = cudf::size_type{2}; + + auto const expected = std::vector(fps + split, fps + size); + auto const col = cudf::test::fixed_point_column_wrapper(reps, reps + size, scale); + auto const splits = std::vector{split}; + auto result = cudf::split(col, splits); + + auto host_data = cudf::test::to_host(result.back()); + + EXPECT_TRUE(std::equal(expected.begin(), expected.end(), host_data.first.begin())); +} + CUDF_TEST_PROGRAM_MAIN()