diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 453707e4559..98f6ed40502 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -210,8 +210,8 @@ add_library(cudf src/groupby/sort/sort_helper.cu src/hash/hashing.cu src/interop/dlpack.cpp - src/interop/from_arrow.cpp - src/interop/to_arrow.cpp + src/interop/from_arrow.cu + src/interop/to_arrow.cu src/io/avro/avro.cpp src/io/avro/avro_gpu.cu src/io/avro/reader_impl.cu diff --git a/cpp/src/interop/from_arrow.cpp b/cpp/src/interop/from_arrow.cu similarity index 88% rename from cpp/src/interop/from_arrow.cpp rename to cpp/src/interop/from_arrow.cu index 99c9b386a15..ee02fadc017 100644 --- a/cpp/src/interop/from_arrow.cpp +++ b/cpp/src/interop/from_arrow.cu @@ -13,12 +13,12 @@ * See the License for the specific language governing permissions and * limitations under the License. */ - #include #include #include #include #include +#include #include #include #include @@ -34,6 +34,8 @@ #include +#include + namespace cudf { namespace detail { @@ -54,7 +56,7 @@ data_type arrow_to_cudf_type(arrow::DataType const& arrow_type) case arrow::Type::DOUBLE: return data_type(type_id::FLOAT64); case arrow::Type::DATE32: return data_type(type_id::TIMESTAMP_DAYS); case arrow::Type::TIMESTAMP: { - arrow::TimestampType const* type = static_cast(&arrow_type); + auto type = static_cast(&arrow_type); switch (type->unit()) { case arrow::TimeUnit::type::SECOND: return data_type(type_id::TIMESTAMP_SECONDS); case arrow::TimeUnit::type::MILLI: return data_type(type_id::TIMESTAMP_MILLISECONDS); @@ -64,7 +66,7 @@ data_type arrow_to_cudf_type(arrow::DataType const& arrow_type) } } case arrow::Type::DURATION: { - arrow::DurationType const* type = static_cast(&arrow_type); + auto type = static_cast(&arrow_type); switch (type->unit()) { case arrow::TimeUnit::type::SECOND: return data_type(type_id::DURATION_SECONDS); case arrow::TimeUnit::type::MILLI: return data_type(type_id::DURATION_MILLISECONDS); @@ -76,6 +78,10 @@ data_type arrow_to_cudf_type(arrow::DataType const& arrow_type) case arrow::Type::STRING: return data_type(type_id::STRING); case arrow::Type::DICTIONARY: return data_type(type_id::DICTIONARY32); 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()}; + } case arrow::Type::STRUCT: return data_type(type_id::STRUCT); default: CUDF_FAIL("Unsupported type_id conversion to cudf"); } @@ -174,6 +180,54 @@ std::unique_ptr get_column(arrow::Array const& array, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr); +template <> +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; + + 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); + + CUDA_TRY(cudaMemcpyAsync( + reinterpret_cast(buf.data()), + reinterpret_cast(data_buffer->address()) + array.offset() * sizeof(DeviceType), + buf.size() * sizeof(DeviceType), + 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); + // If array is sliced, we have to copy whole mask and then take copy. + return (num_rows == static_cast(data_buffer->size() / sizeof(DeviceType))) + ? *temp_mask.release() + : cudf::detail::copy_bitmask(static_cast(temp_mask->data()), + array.offset(), + array.offset() + num_rows, + stream, + mr); + } + return rmm::device_buffer{}; + }(); + + return std::make_unique(type, num_rows, out_buf.release(), std::move(null_mask)); +} + template <> std::unique_ptr dispatch_to_cudf_column::operator()( arrow::Array const& array, diff --git a/cpp/src/interop/to_arrow.cpp b/cpp/src/interop/to_arrow.cu similarity index 88% rename from cpp/src/interop/to_arrow.cpp rename to cpp/src/interop/to_arrow.cu index 4bc50b21718..d9be3316f9d 100644 --- a/cpp/src/interop/to_arrow.cpp +++ b/cpp/src/interop/to_arrow.cu @@ -17,6 +17,7 @@ #include #include #include +#include #include #include #include @@ -30,6 +31,9 @@ #include #include +#include +#include + namespace cudf { namespace detail { namespace { @@ -135,6 +139,49 @@ struct dispatch_to_arrow { } }; +template <> +std::shared_ptr dispatch_to_arrow::operator()( + column_view input, + cudf::type_id id, + column_metadata const& metadata, + arrow::MemoryPool* ar_mr, + rmm::cuda_stream_view stream) +{ + using DeviceType = int64_t; + size_type const BIT_WIDTH_RATIO = 2; // Array::Type:type::DECIMAL (128) / int64_t + + rmm::device_uvector buf(input.size() * BIT_WIDTH_RATIO, stream); + + auto count = thrust::make_counting_iterator(0); + + thrust::for_each(count, + count + input.size(), + [in = input.begin(), out = buf.data()] __device__(auto in_idx) { + auto const out_idx = in_idx * 2; + out[out_idx] = in[in_idx]; + out[out_idx + 1] = in[in_idx] < 0 ? -1 : 0; + }); + + auto const buf_size_in_bytes = buf.size() * sizeof(DeviceType); + auto result = arrow::AllocateBuffer(buf_size_in_bytes, ar_mr); + CUDF_EXPECTS(result.ok(), "Failed to allocate Arrow buffer for data"); + + std::shared_ptr data_buffer = std::move(result.ValueOrDie()); + + 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, 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/from_arrow_test.cpp b/cpp/tests/interop/from_arrow_test.cpp index d79307dcbf6..ae8808ba59d 100644 --- a/cpp/tests/interop/from_arrow_test.cpp +++ b/cpp/tests/interop/from_arrow_test.cpp @@ -32,6 +32,8 @@ #include #include +#include + #include std::unique_ptr get_cudf_table() @@ -76,17 +78,17 @@ TEST_F(FromArrowTest, EmptyTable) TEST_F(FromArrowTest, DateTimeTable) { - auto data = {1, 2, 3, 4, 5, 6}; + auto data = std::vector{1, 2, 3, 4, 5, 6}; - auto col = - cudf::test::fixed_width_column_wrapper(data); + auto col = cudf::test::fixed_width_column_wrapper( + data.begin(), data.end()); cudf::table_view expected_table_view({col}); std::shared_ptr arr; - arrow::TimestampBuilder timestamp_builder(timestamp(arrow::TimeUnit::type::MILLI), + arrow::TimestampBuilder timestamp_builder(arrow::timestamp(arrow::TimeUnit::type::MILLI), arrow::default_memory_pool()); - timestamp_builder.AppendValues(std::vector{1, 2, 3, 4, 5, 6}); + timestamp_builder.AppendValues(data); CUDF_EXPECTS(timestamp_builder.Finish(&arr).ok(), "Failed to build array"); std::vector> schema_vector({arrow::field("a", arr->type())}); @@ -337,10 +339,10 @@ TEST_P(FromArrowTestSlice, SliceTest) auto start = std::get<0>(GetParam()); auto end = std::get<1>(GetParam()); - auto sliced_cudf_table = cudf::slice(cudf_table_view, {start, end})[0]; - cudf::table expected_cudf_table{sliced_cudf_table}; - auto sliced_arrow_table = arrow_table->Slice(start, end - start); - auto got_cudf_table = cudf::from_arrow(*sliced_arrow_table); + auto sliced_cudf_table = cudf::slice(cudf_table_view, {start, end})[0]; + auto expected_cudf_table = cudf::table{sliced_cudf_table}; + auto sliced_arrow_table = arrow_table->Slice(start, end - start); + auto got_cudf_table = cudf::from_arrow(*sliced_arrow_table); // This has been added to take-care of empty string column issue with no children if (got_cudf_table->num_rows() == 0 and expected_cudf_table.num_rows() == 0) { @@ -350,6 +352,131 @@ TEST_P(FromArrowTestSlice, SliceTest) } } +template +using fp_wrapper = cudf::test::fixed_point_column_wrapper; + +TEST_F(FromArrowTest, FixedPointTable) +{ + 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}); + 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 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_cudf_table = cudf::from_arrow(*arrow_table); + + CUDF_TEST_EXPECT_TABLES_EQUAL(expected, got_cudf_table->view()); + } +} + +TEST_F(FromArrowTest, FixedPointTableLarge) +{ + 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 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}); + 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 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_cudf_table = cudf::from_arrow(*arrow_table); + + CUDF_TEST_EXPECT_TABLES_EQUAL(expected, got_cudf_table->view()); + } +} + +TEST_F(FromArrowTest, FixedPointTableNulls) +{ + 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, 0, 0}, {1, 1, 1, 1, 1, 1, 0, 0}, scale_type{i}); + 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 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_cudf_table = cudf::from_arrow(*arrow_table); + + CUDF_TEST_EXPECT_TABLES_EQUAL(expected, got_cudf_table->view()); + } +} + +TEST_F(FromArrowTest, FixedPointTableNullsLarge) +{ + 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 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, transform, scale_type{i}); + 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 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_cudf_table = cudf::from_arrow(*arrow_table); + + CUDF_TEST_EXPECT_TABLES_EQUAL(expected, got_cudf_table->view()); + } +} + INSTANTIATE_TEST_CASE_P(FromArrowTest, FromArrowTestSlice, ::testing::Values(std::make_tuple(0, 10000), diff --git a/cpp/tests/interop/to_arrow_test.cpp b/cpp/tests/interop/to_arrow_test.cpp index 57275433516..00d625175d0 100644 --- a/cpp/tests/interop/to_arrow_test.cpp +++ b/cpp/tests/interop/to_arrow_test.cpp @@ -353,6 +353,134 @@ TEST_F(ToArrowTest, StructColumn) ASSERT_TRUE(expected_arrow_table->Equals(*got_arrow_table, true)); } +template +using fp_wrapper = cudf::test::fixed_point_column_wrapper; + +TEST_F(ToArrowTest, FixedPointTable) +{ + 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}); + + 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 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, FixedPointTableLarge) +{ + 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}); + 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 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)); + + 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, FixedPointTableNullsSimple) +{ + 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, 0, 0}, {1, 1, 1, 1, 1, 1, 0, 0}, scale_type{i}); + 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(); + + 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, FixedPointTableNulls) +{ + 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 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}); + 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(); + } + + 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)); + } +} + struct ToArrowTestSlice : public ToArrowTest, public ::testing::WithParamInterface> { diff --git a/python/cudf/cudf/core/column/column.py b/python/cudf/cudf/core/column/column.py index 6a1600d6461..547e298cc83 100644 --- a/python/cudf/cudf/core/column/column.py +++ b/python/cudf/cudf/core/column/column.py @@ -358,7 +358,7 @@ def to_arrow(self) -> pa.Array: pa.null(), len(self), [pa.py_buffer((b""))] ) - return libcudf.interop.to_arrow( + result = libcudf.interop.to_arrow( libcudf.table.Table( cudf.core.column_accessor.ColumnAccessor({"None": self}) ), @@ -366,6 +366,14 @@ def to_arrow(self) -> pa.Array: keep_index=False, )["None"].chunk(0) + if isinstance(self.dtype, cudf.Decimal64Dtype): + result = result.view( + pa.decimal128( + scale=result.type.scale, precision=self.dtype.precision + ) + ) + return result + @classmethod def from_arrow(cls, array: pa.Array) -> ColumnBase: """ @@ -430,10 +438,14 @@ def from_arrow(cls, array: pa.Array) -> ColumnBase: elif isinstance(array.type, pa.Decimal128Type): return cudf.core.column.DecimalColumn.from_arrow(array) - return libcudf.interop.from_arrow(data, data.column_names)._data[ + result = libcudf.interop.from_arrow(data, data.column_names)._data[ "None" ] + if isinstance(result.dtype, cudf.Decimal64Dtype): + result.dtype.precision = array.type.precision + return result + def _get_mask_as_column(self) -> ColumnBase: return libcudf.transform.mask_to_bools( self.base_mask, self.offset, self.offset + len(self)