diff --git a/cpp/src/interop/from_arrow.cu b/cpp/src/interop/from_arrow.cu index edd3ce2ed07..99b657fb9d5 100644 --- a/cpp/src/interop/from_arrow.cu +++ b/cpp/src/interop/from_arrow.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -79,7 +79,7 @@ data_type arrow_to_cudf_type(arrow::DataType const& arrow_type) case arrow::Type::LIST: return data_type(type_id::LIST); case arrow::Type::DECIMAL: { auto const type = static_cast(&arrow_type); - return data_type{type_id::DECIMAL64, -type->scale()}; + return data_type{type_id::DECIMAL128, -type->scale()}; } case arrow::Type::STRUCT: return data_type(type_id::STRUCT); default: CUDF_FAIL("Unsupported type_id conversion to cudf"); @@ -177,35 +177,27 @@ std::unique_ptr get_column(arrow::Array const& array, rmm::mr::device_memory_resource* mr); template <> -std::unique_ptr dispatch_to_cudf_column::operator()( +std::unique_ptr dispatch_to_cudf_column::operator()( arrow::Array const& array, data_type type, bool skip_mask, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - using DeviceType = int64_t; + using DeviceType = __int128_t; - auto constexpr BIT_WIDTH_RATIO = 2; // Array::Type:type::DECIMAL (128) / int64_t - auto data_buffer = array.data()->buffers[1]; - auto const num_rows = static_cast(array.length()); - - rmm::device_uvector buf(num_rows * BIT_WIDTH_RATIO, stream); - rmm::device_uvector out_buf(num_rows, stream, mr); + auto data_buffer = array.data()->buffers[1]; + auto const num_rows = static_cast(array.length()); + auto col = make_fixed_width_column(type, num_rows, mask_state::UNALLOCATED, stream, mr); + auto mutable_column_view = col->mutable_view(); CUDA_TRY(cudaMemcpyAsync( - reinterpret_cast(buf.data()), + mutable_column_view.data(), reinterpret_cast(data_buffer->address()) + array.offset() * sizeof(DeviceType), - buf.size() * sizeof(DeviceType), + sizeof(DeviceType) * num_rows, cudaMemcpyDefault, stream.value())); - auto every_other = [] __device__(size_type i) { return 2 * i; }; - auto gather_map = cudf::detail::make_counting_transform_iterator(0, every_other); - - thrust::gather( - rmm::exec_policy(stream), gather_map, gather_map + num_rows, buf.data(), out_buf.data()); - auto null_mask = [&] { if (not skip_mask and array.null_bitmap_data()) { auto temp_mask = get_mask_buffer(array, stream, mr); @@ -221,7 +213,8 @@ std::unique_ptr dispatch_to_cudf_column::operator()( return rmm::device_buffer{}; }(); - return std::make_unique(type, num_rows, out_buf.release(), std::move(null_mask)); + col->set_null_mask(std::move(null_mask)); + return col; } template <> diff --git a/cpp/src/interop/to_arrow.cu b/cpp/src/interop/to_arrow.cu index e6db5807dde..27e47061b67 100644 --- a/cpp/src/interop/to_arrow.cu +++ b/cpp/src/interop/to_arrow.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -176,6 +176,40 @@ std::shared_ptr dispatch_to_arrow::operator()( return std::make_shared(data); } +template <> +std::shared_ptr dispatch_to_arrow::operator()( + column_view input, + cudf::type_id, + column_metadata const&, + arrow::MemoryPool* ar_mr, + rmm::cuda_stream_view stream) +{ + using DeviceType = __int128_t; + + rmm::device_uvector buf(input.size(), stream); + + thrust::copy(rmm::exec_policy(stream), // + input.begin(), + input.end(), + buf.begin()); + + auto const buf_size_in_bytes = buf.size() * sizeof(DeviceType); + auto data_buffer = allocate_arrow_buffer(buf_size_in_bytes, ar_mr); + + CUDA_TRY(cudaMemcpyAsync(data_buffer->mutable_data(), + buf.data(), + buf_size_in_bytes, + cudaMemcpyDeviceToHost, + stream.value())); + + auto type = arrow::decimal(18, -input.type().scale()); + auto mask = fetch_mask_buffer(input, ar_mr, stream); + auto buffers = std::vector>{mask, std::move(data_buffer)}; + auto data = std::make_shared(type, input.size(), buffers); + + return std::make_shared(data); +} + template <> std::shared_ptr dispatch_to_arrow::operator()(column_view input, cudf::type_id id, diff --git a/cpp/tests/interop/arrow_utils.hpp b/cpp/tests/interop/arrow_utils.hpp index 78967fbaf30..4da8170e75f 100644 --- a/cpp/tests/interop/arrow_utils.hpp +++ b/cpp/tests/interop/arrow_utils.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -176,3 +176,27 @@ std::shared_ptr get_arrow_list_array( std::pair, std::shared_ptr> get_tables( cudf::size_type length = 10000); + +template +[[nodiscard]] auto make_decimal128_arrow_array(std::vector const& data, + std::optional> const& validity, + int32_t scale) -> std::shared_ptr +{ + auto constexpr BIT_WIDTH_RATIO = sizeof(__int128_t) / sizeof(T); + + std::shared_ptr arr; + arrow::Decimal128Builder decimal_builder(arrow::decimal(18, -scale), + arrow::default_memory_pool()); + + for (T i = 0; i < static_cast(data.size() / BIT_WIDTH_RATIO); ++i) { + if (validity.has_value() and not validity.value()[i]) { + decimal_builder.AppendNull(); + } else { + decimal_builder.Append(reinterpret_cast(data.data() + BIT_WIDTH_RATIO * i)); + } + } + + CUDF_EXPECTS(decimal_builder.Finish(&arr).ok(), "Failed to build array"); + + return arr; +} diff --git a/cpp/tests/interop/from_arrow_test.cpp b/cpp/tests/interop/from_arrow_test.cpp index 946ac7fc891..f5176cda768 100644 --- a/cpp/tests/interop/from_arrow_test.cpp +++ b/cpp/tests/interop/from_arrow_test.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -354,21 +354,16 @@ TEST_P(FromArrowTestSlice, SliceTest) template using fp_wrapper = cudf::test::fixed_point_column_wrapper; -TEST_F(FromArrowTest, FixedPointTable) +TEST_F(FromArrowTest, FixedPoint128Table) { using namespace numeric; - auto constexpr BIT_WIDTH_RATIO = 2; // Array::Type:type::DECIMAL (128) / int64_t - for (auto const i : {3, 2, 1, 0, -1, -2, -3}) { - auto const data = std::vector{1, 0, 2, 0, 3, 0, 4, 0, 5, 0, 6, 0}; - auto const col = fp_wrapper({1, 2, 3, 4, 5, 6}, scale_type{i}); + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { + auto const data = std::vector<__int128_t>{1, 2, 3, 4, 5, 6}; + auto const col = fp_wrapper<__int128_t>(data.cbegin(), data.cend(), scale_type{scale}); auto const expected = cudf::table_view({col}); - std::shared_ptr arr; - arrow::Decimal128Builder decimal_builder(arrow::decimal(10, -i), arrow::default_memory_pool()); - decimal_builder.AppendValues(reinterpret_cast(data.data()), - data.size() / BIT_WIDTH_RATIO); - CUDF_EXPECTS(decimal_builder.Finish(&arr).ok(), "Failed to build array"); + auto const arr = make_decimal128_arrow_array(data, std::nullopt, scale); auto const field = arrow::field("a", arr->type()); auto const schema_vector = std::vector>({field}); @@ -381,24 +376,18 @@ TEST_F(FromArrowTest, FixedPointTable) } } -TEST_F(FromArrowTest, FixedPointTableLarge) +TEST_F(FromArrowTest, FixedPoint128TableLarge) { using namespace numeric; - auto constexpr BIT_WIDTH_RATIO = 2; // Array::Type:type::DECIMAL (128) / int64_t - auto constexpr NUM_ELEMENTS = 1000; + auto constexpr NUM_ELEMENTS = 1000; - for (auto const i : {3, 2, 1, 0, -1, -2, -3}) { - auto every_other = [](auto i) { return i % BIT_WIDTH_RATIO ? 0 : i / BIT_WIDTH_RATIO; }; - auto transform = cudf::detail::make_counting_transform_iterator(BIT_WIDTH_RATIO, every_other); - auto const data = std::vector(transform, transform + NUM_ELEMENTS * BIT_WIDTH_RATIO); - auto iota = thrust::make_counting_iterator(1); - auto const col = fp_wrapper(iota, iota + NUM_ELEMENTS, scale_type{i}); + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { + auto iota = thrust::make_counting_iterator(1); + auto const data = std::vector<__int128_t>(iota, iota + NUM_ELEMENTS); + auto const col = fp_wrapper<__int128_t>(iota, iota + NUM_ELEMENTS, scale_type{scale}); auto const expected = cudf::table_view({col}); - std::shared_ptr arr; - arrow::Decimal128Builder decimal_builder(arrow::decimal(10, -i), arrow::default_memory_pool()); - decimal_builder.AppendValues(reinterpret_cast(data.data()), NUM_ELEMENTS); - CUDF_EXPECTS(decimal_builder.Finish(&arr).ok(), "Failed to build array"); + auto const arr = make_decimal128_arrow_array(data, std::nullopt, scale); auto const field = arrow::field("a", arr->type()); auto const schema_vector = std::vector>({field}); @@ -411,25 +400,18 @@ TEST_F(FromArrowTest, FixedPointTableLarge) } } -TEST_F(FromArrowTest, FixedPointTableNulls) +TEST_F(FromArrowTest, FixedPoint128TableNulls) { using namespace numeric; - auto constexpr BIT_WIDTH_RATIO = 2; // Array::Type:type::DECIMAL (128) / int64_t - for (auto const i : {3, 2, 1, 0, -1, -2, -3}) { - auto const data = std::vector{1, 0, 2, 0, 3, 0, 4, 0, 5, 0, 6, 0}; + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { + auto const data = std::vector<__int128_t>{1, 2, 3, 4, 5, 6, 0, 0}; + auto const validity = std::vector{1, 1, 1, 1, 1, 1, 0, 0}; auto const col = - fp_wrapper({1, 2, 3, 4, 5, 6, 0, 0}, {1, 1, 1, 1, 1, 1, 0, 0}, scale_type{i}); + fp_wrapper<__int128_t>({1, 2, 3, 4, 5, 6, 0, 0}, {1, 1, 1, 1, 1, 1, 0, 0}, scale_type{scale}); auto const expected = cudf::table_view({col}); - std::shared_ptr arr; - arrow::Decimal128Builder decimal_builder(arrow::decimal(10, -i), arrow::default_memory_pool()); - decimal_builder.AppendValues(reinterpret_cast(data.data()), - data.size() / BIT_WIDTH_RATIO); - decimal_builder.AppendNull(); - decimal_builder.AppendNull(); - - CUDF_EXPECTS(decimal_builder.Finish(&arr).ok(), "Failed to build array"); + auto const arr = make_decimal128_arrow_array(data, validity, scale); auto const field = arrow::field("a", arr->type()); auto const schema_vector = std::vector>({field}); @@ -442,28 +424,21 @@ TEST_F(FromArrowTest, FixedPointTableNulls) } } -TEST_F(FromArrowTest, FixedPointTableNullsLarge) +TEST_F(FromArrowTest, FixedPoint128TableNullsLarge) { using namespace numeric; - auto constexpr BIT_WIDTH_RATIO = 2; // Array::Type:type::DECIMAL (128) / int64_t - auto constexpr NUM_ELEMENTS = 1000; + auto constexpr NUM_ELEMENTS = 1000; - for (auto const i : {3, 2, 1, 0, -1, -2, -3}) { - auto every_other = [](auto i) { return i % BIT_WIDTH_RATIO ? 0 : i / BIT_WIDTH_RATIO; }; - auto transform = cudf::detail::make_counting_transform_iterator(BIT_WIDTH_RATIO, every_other); - auto const data = std::vector(transform, transform + NUM_ELEMENTS * BIT_WIDTH_RATIO); + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { + auto every_other = [](auto i) { return i % 2 ? 0 : 1; }; + auto validity = cudf::detail::make_counting_transform_iterator(0, every_other); auto iota = thrust::make_counting_iterator(1); - auto const col = fp_wrapper(iota, iota + NUM_ELEMENTS, transform, scale_type{i}); + auto const data = std::vector<__int128_t>(iota, iota + NUM_ELEMENTS); + auto const col = fp_wrapper<__int128_t>(iota, iota + NUM_ELEMENTS, validity, scale_type{scale}); auto const expected = cudf::table_view({col}); - std::shared_ptr arr; - arrow::Decimal128Builder decimal_builder(arrow::decimal(10, -i), arrow::default_memory_pool()); - for (int64_t i = 0; i < NUM_ELEMENTS / BIT_WIDTH_RATIO; ++i) { - decimal_builder.Append(reinterpret_cast(data.data() + 4 * i)); - decimal_builder.AppendNull(); - } - - CUDF_EXPECTS(decimal_builder.Finish(&arr).ok(), "Failed to build array"); + auto const arr = make_decimal128_arrow_array( + data, std::vector(validity, validity + NUM_ELEMENTS), scale); auto const field = arrow::field("a", arr->type()); auto const schema_vector = std::vector>({field}); diff --git a/cpp/tests/interop/to_arrow_test.cpp b/cpp/tests/interop/to_arrow_test.cpp index 98031f42a9c..52f2d5709d2 100644 --- a/cpp/tests/interop/to_arrow_test.cpp +++ b/cpp/tests/interop/to_arrow_test.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -356,21 +356,16 @@ TEST_F(ToArrowTest, StructColumn) template using fp_wrapper = cudf::test::fixed_point_column_wrapper; -TEST_F(ToArrowTest, FixedPointTable) +TEST_F(ToArrowTest, FixedPoint64Table) { using namespace numeric; - auto constexpr const BIT_WIDTH_RATIO = 2; // Array::Type:type::DECIMAL (128) / int64_t - - for (auto const i : {3, 2, 1, 0, -1, -2, -3}) { - auto const col = fp_wrapper({-1, 2, 3, 4, 5, 6}, scale_type{i}); - auto const input = cudf::table_view({col}); + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { + auto const col = fp_wrapper({-1, 2, 3, 4, 5, 6}, scale_type{scale}); + auto const input = cudf::table_view({col}); auto const expect_data = std::vector{-1, -1, 2, 0, 3, 0, 4, 0, 5, 0, 6, 0}; - std::shared_ptr arr; - arrow::Decimal128Builder decimal_builder(arrow::decimal(18, -i), arrow::default_memory_pool()); - decimal_builder.AppendValues(reinterpret_cast(expect_data.data()), - expect_data.size() / BIT_WIDTH_RATIO); - CUDF_EXPECTS(decimal_builder.Finish(&arr).ok(), "Failed to build array"); + + auto const arr = make_decimal128_arrow_array(expect_data, std::nullopt, scale); auto const field = arrow::field("a", arr->type()); auto const schema_vector = std::vector>({field}); @@ -383,29 +378,69 @@ TEST_F(ToArrowTest, FixedPointTable) } } -TEST_F(ToArrowTest, FixedPointTableLarge) +TEST_F(ToArrowTest, FixedPoint128Table) +{ + using namespace numeric; + + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { + auto const col = fp_wrapper<__int128_t>({-1, 2, 3, 4, 5, 6}, scale_type{scale}); + auto const input = cudf::table_view({col}); + auto const expect_data = std::vector<__int128_t>{-1, 2, 3, 4, 5, 6}; + + auto const arr = make_decimal128_arrow_array(expect_data, std::nullopt, scale); + + auto const field = arrow::field("a", arr->type()); + auto const schema_vector = std::vector>({field}); + auto const schema = std::make_shared(schema_vector); + auto const expected_arrow_table = arrow::Table::Make(schema, {arr}); + + auto got_arrow_table = cudf::to_arrow(input, {{"a"}}); + + ASSERT_TRUE(expected_arrow_table->Equals(*got_arrow_table, true)); + } +} + +TEST_F(ToArrowTest, FixedPoint64TableLarge) { using namespace numeric; auto constexpr BIT_WIDTH_RATIO = 2; // Array::Type:type::DECIMAL (128) / int64_t auto constexpr NUM_ELEMENTS = 1000; - for (auto const i : {3, 2, 1, 0, -1, -2, -3}) { - auto iota = thrust::make_counting_iterator(1); - auto const col = fp_wrapper(iota, iota + NUM_ELEMENTS, scale_type{i}); + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { + auto const iota = thrust::make_counting_iterator(1); + auto const col = fp_wrapper(iota, iota + NUM_ELEMENTS, scale_type{scale}); auto const input = cudf::table_view({col}); - auto every_other = [](auto i) { return i % 2 == 0 ? i / 2 : 0; }; - auto transform = cudf::detail::make_counting_transform_iterator(2, every_other); + auto const every_other = [](auto i) { return i % 2 == 0 ? i / 2 : 0; }; + auto const transform = cudf::detail::make_counting_transform_iterator(2, every_other); auto const expect_data = std::vector{transform, transform + NUM_ELEMENTS * BIT_WIDTH_RATIO}; - std::shared_ptr arr; - arrow::Decimal128Builder decimal_builder(arrow::decimal(18, -i), arrow::default_memory_pool()); - // Note: For some reason, decimal_builder.AppendValues with NUM_ELEMENTS >= 1000 doesn't work - for (int i = 0; i < NUM_ELEMENTS; ++i) - decimal_builder.Append(reinterpret_cast(expect_data.data() + 2 * i)); + auto const arr = make_decimal128_arrow_array(expect_data, std::nullopt, scale); - CUDF_EXPECTS(decimal_builder.Finish(&arr).ok(), "Failed to build array"); + auto const field = arrow::field("a", arr->type()); + auto const schema_vector = std::vector>({field}); + auto const schema = std::make_shared(schema_vector); + auto const expected_arrow_table = arrow::Table::Make(schema, {arr}); + + auto got_arrow_table = cudf::to_arrow(input, {{"a"}}); + + ASSERT_TRUE(expected_arrow_table->Equals(*got_arrow_table, true)); + } +} + +TEST_F(ToArrowTest, FixedPoint128TableLarge) +{ + using namespace numeric; + auto constexpr NUM_ELEMENTS = 1000; + + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { + auto const iota = thrust::make_counting_iterator(1); + auto const col = fp_wrapper<__int128_t>(iota, iota + NUM_ELEMENTS, scale_type{scale}); + auto const input = cudf::table_view({col}); + auto const expect_data = std::vector<__int128_t>{iota, iota + NUM_ELEMENTS}; + + auto const arr = make_decimal128_arrow_array(expect_data, std::nullopt, scale); auto const field = arrow::field("a", arr->type()); auto const schema_vector = std::vector>({field}); @@ -418,25 +453,42 @@ TEST_F(ToArrowTest, FixedPointTableLarge) } } -TEST_F(ToArrowTest, FixedPointTableNullsSimple) +TEST_F(ToArrowTest, FixedPoint64TableNullsSimple) { using namespace numeric; - auto constexpr BIT_WIDTH_RATIO = 2; // Array::Type:type::DECIMAL (128) / int64_t - for (auto const i : {3, 2, 1, 0, -1, -2, -3}) { - auto const data = std::vector{1, 0, 2, 0, 3, 0, 4, 0, 5, 0, 6, 0}; + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { + auto const data = std::vector{1, 0, 2, 0, 3, 0, 4, 0, 5, 0, 6, 0, 0, 0, 0, 0}; + auto const validity = std::vector{1, 1, 1, 1, 1, 1, 0, 0}; auto const col = - fp_wrapper({1, 2, 3, 4, 5, 6, 0, 0}, {1, 1, 1, 1, 1, 1, 0, 0}, scale_type{i}); + fp_wrapper({1, 2, 3, 4, 5, 6, 0, 0}, {1, 1, 1, 1, 1, 1, 0, 0}, scale_type{scale}); auto const input = cudf::table_view({col}); - std::shared_ptr arr; - arrow::Decimal128Builder decimal_builder(arrow::decimal(18, -i), arrow::default_memory_pool()); - decimal_builder.AppendValues(reinterpret_cast(data.data()), - data.size() / BIT_WIDTH_RATIO); - decimal_builder.AppendNull(); - decimal_builder.AppendNull(); + auto const arr = make_decimal128_arrow_array(data, validity, scale); - CUDF_EXPECTS(decimal_builder.Finish(&arr).ok(), "Failed to build array"); + auto const field = arrow::field("a", arr->type()); + auto const schema_vector = std::vector>({field}); + auto const schema = std::make_shared(schema_vector); + auto const arrow_table = arrow::Table::Make(schema, {arr}); + + auto got_arrow_table = cudf::to_arrow(input, {{"a"}}); + + ASSERT_TRUE(arrow_table->Equals(*got_arrow_table, true)); + } +} + +TEST_F(ToArrowTest, FixedPoint128TableNullsSimple) +{ + using namespace numeric; + + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { + auto const data = std::vector<__int128_t>{1, 2, 3, 4, 5, 6, 0, 0}; + auto const validity = std::vector{1, 1, 1, 1, 1, 1, 0, 0}; + auto const col = + fp_wrapper<__int128_t>({1, 2, 3, 4, 5, 6, 0, 0}, {1, 1, 1, 1, 1, 1, 0, 0}, scale_type{scale}); + auto const input = cudf::table_view({col}); + + auto const arr = make_decimal128_arrow_array(data, validity, scale); auto const field = arrow::field("a", arr->type()); auto const schema_vector = std::vector>({field}); @@ -449,26 +501,20 @@ TEST_F(ToArrowTest, FixedPointTableNullsSimple) } } -TEST_F(ToArrowTest, FixedPointTableNulls) +TEST_F(ToArrowTest, FixedPoint64TableNulls) { using namespace numeric; - auto constexpr BIT_WIDTH_RATIO = 2; // Array::Type:type::DECIMAL (128) / int64_t - for (auto const i : {3, 2, 1, 0, -1, -2, -3}) { + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { auto const col = fp_wrapper( - {1, 2, 3, 4, 5, 6, 7, 8, 9, 10}, {1, 0, 1, 0, 1, 0, 1, 0, 1, 0}, scale_type{i}); + {1, 2, 3, 4, 5, 6, 7, 8, 9, 10}, {1, 0, 1, 0, 1, 0, 1, 0, 1, 0}, scale_type{scale}); auto const input = cudf::table_view({col}); auto const expect_data = std::vector{1, 0, 2, 0, 3, 0, 4, 0, 5, 0, 6, 0, 7, 0, 8, 0, 9, 0, 10, 0}; - std::shared_ptr arr; - arrow::Decimal128Builder decimal_builder(arrow::decimal(18, -i), arrow::default_memory_pool()); - for (int64_t i = 0; i < input.column(0).size() / BIT_WIDTH_RATIO; ++i) { - decimal_builder.Append(reinterpret_cast(expect_data.data() + 4 * i)); - decimal_builder.AppendNull(); - } + auto const validity = std::vector{1, 0, 1, 0, 1, 0, 1, 0, 1, 0}; - CUDF_EXPECTS(decimal_builder.Finish(&arr).ok(), "Failed to build array"); + auto arr = make_decimal128_arrow_array(expect_data, validity, scale); auto const field = arrow::field("a", arr->type()); auto const schema_vector = std::vector>({field}); @@ -481,6 +527,31 @@ TEST_F(ToArrowTest, FixedPointTableNulls) } } +TEST_F(ToArrowTest, FixedPoint128TableNulls) +{ + using namespace numeric; + + for (auto const scale : {3, 2, 1, 0, -1, -2, -3}) { + auto const col = fp_wrapper<__int128_t>( + {1, 2, 3, 4, 5, 6, 7, 8, 9, 10}, {1, 0, 1, 0, 1, 0, 1, 0, 1, 0}, scale_type{scale}); + auto const input = cudf::table_view({col}); + + auto const expect_data = std::vector<__int128_t>{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; + auto const validity = std::vector{1, 0, 1, 0, 1, 0, 1, 0, 1, 0}; + + auto arr = make_decimal128_arrow_array(expect_data, validity, scale); + + auto const field = arrow::field("a", arr->type()); + auto const schema_vector = std::vector>({field}); + auto const schema = std::make_shared(schema_vector); + auto const expected_arrow_table = arrow::Table::Make(schema, {arr}); + + auto const got_arrow_table = cudf::to_arrow(input, {{"a"}}); + + ASSERT_TRUE(expected_arrow_table->Equals(*got_arrow_table, true)); + } +} + struct ToArrowTestSlice : public ToArrowTest, public ::testing::WithParamInterface> {