Implement cudf::copy_range for decimal32 and decimal64 (#6843)

CHANGELOG.md

@@ -11,6 +11,8 @@
 - PR #6777 Implement `cudf::unary_operation` for `decimal32` & `decimal64`
 - PR #6729 Implement `cudf::cast` for `decimal32/64` to/from different `type_id`
 - PR #6792 Implement `cudf::clamp` for `decimal32` and `decimal64`
+- PR #6805 Implement `cudf::detail::copy_if` for `decimal32` and `decimal64`
+- PR #6843 Implement `cudf::copy_range` for `decimal32` and `decimal64`
 - PR #6528 Enable `fixed_point` binary operations
 - PR #6460 Add is_timestamp format check API
 - PR #6568 Add function to create hashed vocabulary file from raw vocabulary
@@ -29,7 +31,6 @@
 - PR #6768 Add support for scatter() on list columns
 - PR #6796 Add create_metadata_file in dask_cudf
 - PR #6765 Cupy fallback for __array_function__ and __array_ufunc__ for cudf.Series
-- PR #6805 Implement `cudf::detail::copy_if` for `decimal32` and `decimal64`
 
 ## Improvements
 

cpp/include/cudf/detail/copy_range.cuh

@@ -144,7 +144,7 @@ void copy_range(SourceValueIterator source_value_begin,
   using T = typename std::iterator_traits<SourceValueIterator>::value_type;
 
   // this code assumes that source and target have the same type.
-  CUDF_EXPECTS(type_to_id<T>() == target.type().id(), "the data type mismatch");
+  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(target.type().id()), "data type mismatch");
 
   auto warp_aligned_begin_lower_bound = cudf::util::round_down_safe(target_begin, warp_size);
   auto warp_aligned_end_upper_bound   = cudf::util::round_up_safe(target_end, warp_size);

cpp/src/copying/copy_range.cu

@@ -108,8 +108,9 @@ struct out_of_place_copy_range_dispatch {
     }
 
     if (source_end != source_begin) {  // otherwise no-op
+      using Type    = cudf::device_storage_type_t<T>;
       auto ret_view = p_ret->mutable_view();
-      in_place_copy_range<T>(source, ret_view, source_begin, source_end, target_begin, stream);
+      in_place_copy_range<Type>(source, ret_view, source_begin, source_end, target_begin, stream);
     }
 
     return p_ret;
@@ -149,28 +150,6 @@ std::unique_ptr<cudf::column> out_of_place_copy_range_dispatch::operator()<cudf:
   }
 }
 
-template <>
-std::unique_ptr<cudf::column> out_of_place_copy_range_dispatch::operator()<numeric::decimal64>(
-  cudf::size_type source_begin,
-  cudf::size_type source_end,
-  cudf::size_type target_begin,
-  rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
-{
-  CUDF_FAIL("decimal64 type not supported");
-}
-
-template <>
-std::unique_ptr<cudf::column> out_of_place_copy_range_dispatch::operator()<numeric::decimal32>(
-  cudf::size_type source_begin,
-  cudf::size_type source_end,
-  cudf::size_type target_begin,
-  rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
-{
-  CUDF_FAIL("decimal32 type not supported");
-}
-
 template <>
 std::unique_ptr<cudf::column> out_of_place_copy_range_dispatch::operator()<cudf::dictionary32>(
   cudf::size_type source_begin,

cpp/tests/copying/copy_range_tests.cpp

@@ -25,6 +25,7 @@
 #include <cudf/copying.hpp>
 #include <cudf/dictionary/encode.hpp>
 
+#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
 
 auto all_valid  = [](cudf::size_type row) { return true; };
@@ -463,3 +464,57 @@ TEST_F(CopyRangeErrorTestFixture, DTypeMismatch)
   EXPECT_THROW(cudf::copy_range(dict_source->view(), dict_target->view(), 0, 100, 0),
                cudf::logic_error);
 }
+
+template <typename T>
+struct FixedPointTypes : public cudf::test::BaseFixture {
+};
+
+TYPED_TEST_CASE(FixedPointTypes, cudf::test::FixedPointTypes);
+
+TYPED_TEST(FixedPointTypes, FixedPointSimple)
+{
+  using namespace numeric;
+  using decimalXX  = TypeParam;
+  using RepType    = cudf::device_storage_type_t<decimalXX>;
+  using fp_wrapper = cudf::test::fixed_point_column_wrapper<RepType>;
+
+  auto const source   = fp_wrapper{{110, 220, 330, 440, 550, 660}, scale_type{-2}};
+  auto const target   = fp_wrapper{{0, 0, 0, 0, 0, 0}, scale_type{-2}};
+  auto const expected = fp_wrapper{{0, 220, 330, 440, 0, 0}, scale_type{-2}};
+  auto const result   = cudf::copy_range(source, target, 1, 4, 1);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
+TYPED_TEST(FixedPointTypes, FixedPointLarge)
+{
+  using namespace numeric;
+  using namespace cudf::test;
+  using decimalXX  = TypeParam;
+  using RepType    = cudf::device_storage_type_t<decimalXX>;
+  using fp_wrapper = cudf::test::fixed_point_column_wrapper<RepType>;
+
+  auto s = thrust::make_counting_iterator(-1000);
+  auto t = thrust::make_constant_iterator(0);
+  auto e = make_counting_transform_iterator(500, [](int i) { return i < 1000 ? i : 0; });
+
+  auto const source   = fp_wrapper{s, s + 2000, scale_type{-1}};
+  auto const target   = fp_wrapper{t, t + 2000, scale_type{-1}};
+  auto const expected = fp_wrapper{e, e + 2000, scale_type{-1}};
+  auto const result   = cudf::copy_range(source, target, 1500, 2000, 0);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
+TYPED_TEST(FixedPointTypes, FixedPointScaleMismatch)
+{
+  using namespace numeric;
+  using decimalXX  = TypeParam;
+  using RepType    = cudf::device_storage_type_t<decimalXX>;
+  using fp_wrapper = cudf::test::fixed_point_column_wrapper<RepType>;
+
+  auto const source = fp_wrapper{{110, 220, 330, 440, 550, 660}, scale_type{-2}};
+  auto const target = fp_wrapper{{0, 0, 0, 0, 0, 0}, scale_type{-3}};
+
+  EXPECT_THROW(cudf::copy_range(source, target, 1, 4, 1), cudf::logic_error);
+}